mirror of
https://github.com/capstone-engine/llvm-capstone.git
synced 2025-02-25 04:51:38 +00:00
HIP: Directly call ldexp builtins
This commit is contained in:
Matt Arsenault
parent
f14c8eb595
commit
c448e1dabe
@ -304,7 +304,7 @@ float jnf(int __n, float __x) { // TODO: we could use Ahmes multiplication
|
||||
}
|
||||
|
||||
__DEVICE__
|
||||
float ldexpf(float __x, int __e) { return __ocml_ldexp_f32(__x, __e); }
|
||||
float ldexpf(float __x, int __e) { return __builtin_amdgcn_ldexpf(__x, __e); }
|
||||
|
||||
__DEVICE__
|
||||
float lgammaf(float __x) { return __ocml_lgamma_f32(__x); }
|
||||
@ -468,12 +468,12 @@ float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); }
|
||||
|
||||
__DEVICE__
|
||||
float scalblnf(float __x, long int __n) {
|
||||
return (__n < INT_MAX) ? __ocml_scalbn_f32(__x, __n)
|
||||
return (__n < INT_MAX) ? __builtin_amdgcn_ldexpf(__x, __n)
|
||||
: __ocml_scalb_f32(__x, __n);
|
||||
}
|
||||
|
||||
__DEVICE__
|
||||
float scalbnf(float __x, int __n) { return __ocml_scalbn_f32(__x, __n); }
|
||||
float scalbnf(float __x, int __n) { return __builtin_amdgcn_ldexpf(__x, __n); }
|
||||
|
||||
__DEVICE__
|
||||
__RETURN_TYPE __signbitf(float __x) { return __builtin_signbitf(__x); }
|
||||
@ -853,7 +853,7 @@ double jn(int __n, double __x) { // TODO: we could use Ahmes multiplication
|
||||
}
|
||||
|
||||
__DEVICE__
|
||||
double ldexp(double __x, int __e) { return __ocml_ldexp_f64(__x, __e); }
|
||||
double ldexp(double __x, int __e) { return __builtin_amdgcn_ldexp(__x, __e); }
|
||||
|
||||
__DEVICE__
|
||||
double lgamma(double __x) { return __ocml_lgamma_f64(__x); }
|
||||
@ -1025,11 +1025,11 @@ double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); }
|
||||
|
||||
__DEVICE__
|
||||
double scalbln(double __x, long int __n) {
|
||||
return (__n < INT_MAX) ? __ocml_scalbn_f64(__x, __n)
|
||||
return (__n < INT_MAX) ? __builtin_amdgcn_ldexp(__x, __n)
|
||||
: __ocml_scalb_f64(__x, __n);
|
||||
}
|
||||
__DEVICE__
|
||||
double scalbn(double __x, int __n) { return __ocml_scalbn_f64(__x, __n); }
|
||||
double scalbn(double __x, int __n) { return __builtin_amdgcn_ldexp(__x, __n); }
|
||||
|
||||
__DEVICE__
|
||||
__RETURN_TYPE __signbit(double __x) { return __builtin_signbit(__x); }
|
||||
|
||||
@ -1439,13 +1439,13 @@ extern "C" __device__ double test_jn(int x, double y) {
|
||||
|
||||
// DEFAULT-LABEL: @test_ldexpf(
|
||||
// DEFAULT-NEXT: entry:
|
||||
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
|
||||
// DEFAULT-NEXT: ret float [[CALL_I]]
|
||||
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
|
||||
// DEFAULT-NEXT: ret float [[TMP0]]
|
||||
//
|
||||
// FINITEONLY-LABEL: @test_ldexpf(
|
||||
// FINITEONLY-NEXT: entry:
|
||||
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ldexp_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
|
||||
// FINITEONLY-NEXT: ret float [[CALL_I]]
|
||||
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
|
||||
// FINITEONLY-NEXT: ret float [[TMP0]]
|
||||
//
|
||||
extern "C" __device__ float test_ldexpf(float x, int y) {
|
||||
return ldexpf(x, y);
|
||||
@ -1453,13 +1453,13 @@ extern "C" __device__ float test_ldexpf(float x, int y) {
|
||||
|
||||
// DEFAULT-LABEL: @test_ldexp(
|
||||
// DEFAULT-NEXT: entry:
|
||||
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
|
||||
// DEFAULT-NEXT: ret double [[CALL_I]]
|
||||
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
|
||||
// DEFAULT-NEXT: ret double [[TMP0]]
|
||||
//
|
||||
// FINITEONLY-LABEL: @test_ldexp(
|
||||
// FINITEONLY-NEXT: entry:
|
||||
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ldexp_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
|
||||
// FINITEONLY-NEXT: ret double [[CALL_I]]
|
||||
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
|
||||
// FINITEONLY-NEXT: ret double [[TMP0]]
|
||||
//
|
||||
extern "C" __device__ double test_ldexp(double x, int y) {
|
||||
return ldexp(x, y);
|
||||
@ -2688,13 +2688,13 @@ extern "C" __device__ double test_rsqrt(double x) {
|
||||
// DEFAULT-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
|
||||
// DEFAULT: cond.true.i:
|
||||
// DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
|
||||
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
|
||||
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[CONV_I]])
|
||||
// DEFAULT-NEXT: br label [[_ZL8SCALBLNFFL_EXIT:%.*]]
|
||||
// DEFAULT: cond.false.i:
|
||||
// DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR14]]
|
||||
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR14]]
|
||||
// DEFAULT-NEXT: br label [[_ZL8SCALBLNFFL_EXIT]]
|
||||
// DEFAULT: _ZL8scalblnffl.exit:
|
||||
// DEFAULT-NEXT: [[COND_I:%.*]] = phi contract float [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
|
||||
// DEFAULT-NEXT: [[COND_I:%.*]] = phi contract float [ [[TMP0]], [[COND_TRUE_I]] ], [ [[CALL_I]], [[COND_FALSE_I]] ]
|
||||
// DEFAULT-NEXT: ret float [[COND_I]]
|
||||
//
|
||||
// FINITEONLY-LABEL: @test_scalblnf(
|
||||
@ -2703,13 +2703,13 @@ extern "C" __device__ double test_rsqrt(double x) {
|
||||
// FINITEONLY-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
|
||||
// FINITEONLY: cond.true.i:
|
||||
// FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
|
||||
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_scalbn_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
|
||||
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[CONV_I]])
|
||||
// FINITEONLY-NEXT: br label [[_ZL8SCALBLNFFL_EXIT:%.*]]
|
||||
// FINITEONLY: cond.false.i:
|
||||
// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_scalb_f32(float noundef nofpclass(nan inf) [[X]], float noundef nofpclass(nan inf) 0x43E0000000000000) #[[ATTR14]]
|
||||
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_scalb_f32(float noundef nofpclass(nan inf) [[X]], float noundef nofpclass(nan inf) 0x43E0000000000000) #[[ATTR14]]
|
||||
// FINITEONLY-NEXT: br label [[_ZL8SCALBLNFFL_EXIT]]
|
||||
// FINITEONLY: _ZL8scalblnffl.exit:
|
||||
// FINITEONLY-NEXT: [[COND_I:%.*]] = phi nnan ninf contract float [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
|
||||
// FINITEONLY-NEXT: [[COND_I:%.*]] = phi nnan ninf contract float [ [[TMP0]], [[COND_TRUE_I]] ], [ [[CALL_I]], [[COND_FALSE_I]] ]
|
||||
// FINITEONLY-NEXT: ret float [[COND_I]]
|
||||
//
|
||||
extern "C" __device__ float test_scalblnf(float x, long int y) {
|
||||
@ -2722,13 +2722,13 @@ extern "C" __device__ float test_scalblnf(float x, long int y) {
|
||||
// DEFAULT-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
|
||||
// DEFAULT: cond.true.i:
|
||||
// DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
|
||||
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
|
||||
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[CONV_I]])
|
||||
// DEFAULT-NEXT: br label [[_ZL7SCALBLNDL_EXIT:%.*]]
|
||||
// DEFAULT: cond.false.i:
|
||||
// DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR14]]
|
||||
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR14]]
|
||||
// DEFAULT-NEXT: br label [[_ZL7SCALBLNDL_EXIT]]
|
||||
// DEFAULT: _ZL7scalblndl.exit:
|
||||
// DEFAULT-NEXT: [[COND_I:%.*]] = phi contract double [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
|
||||
// DEFAULT-NEXT: [[COND_I:%.*]] = phi contract double [ [[TMP0]], [[COND_TRUE_I]] ], [ [[CALL_I]], [[COND_FALSE_I]] ]
|
||||
// DEFAULT-NEXT: ret double [[COND_I]]
|
||||
//
|
||||
// FINITEONLY-LABEL: @test_scalbln(
|
||||
@ -2737,13 +2737,13 @@ extern "C" __device__ float test_scalblnf(float x, long int y) {
|
||||
// FINITEONLY-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
|
||||
// FINITEONLY: cond.true.i:
|
||||
// FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
|
||||
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_scalbn_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
|
||||
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[CONV_I]])
|
||||
// FINITEONLY-NEXT: br label [[_ZL7SCALBLNDL_EXIT:%.*]]
|
||||
// FINITEONLY: cond.false.i:
|
||||
// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_scalb_f64(double noundef nofpclass(nan inf) [[X]], double noundef nofpclass(nan inf) 0x43E0000000000000) #[[ATTR14]]
|
||||
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_scalb_f64(double noundef nofpclass(nan inf) [[X]], double noundef nofpclass(nan inf) 0x43E0000000000000) #[[ATTR14]]
|
||||
// FINITEONLY-NEXT: br label [[_ZL7SCALBLNDL_EXIT]]
|
||||
// FINITEONLY: _ZL7scalblndl.exit:
|
||||
// FINITEONLY-NEXT: [[COND_I:%.*]] = phi nnan ninf contract double [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
|
||||
// FINITEONLY-NEXT: [[COND_I:%.*]] = phi nnan ninf contract double [ [[TMP0]], [[COND_TRUE_I]] ], [ [[CALL_I]], [[COND_FALSE_I]] ]
|
||||
// FINITEONLY-NEXT: ret double [[COND_I]]
|
||||
//
|
||||
extern "C" __device__ double test_scalbln(double x, long int y) {
|
||||
@ -2752,13 +2752,13 @@ extern "C" __device__ double test_scalbln(double x, long int y) {
|
||||
|
||||
// DEFAULT-LABEL: @test_scalbnf(
|
||||
// DEFAULT-NEXT: entry:
|
||||
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
|
||||
// DEFAULT-NEXT: ret float [[CALL_I]]
|
||||
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
|
||||
// DEFAULT-NEXT: ret float [[TMP0]]
|
||||
//
|
||||
// FINITEONLY-LABEL: @test_scalbnf(
|
||||
// FINITEONLY-NEXT: entry:
|
||||
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_scalbn_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
|
||||
// FINITEONLY-NEXT: ret float [[CALL_I]]
|
||||
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
|
||||
// FINITEONLY-NEXT: ret float [[TMP0]]
|
||||
//
|
||||
extern "C" __device__ float test_scalbnf(float x, int y) {
|
||||
return scalbnf(x, y);
|
||||
@ -2766,13 +2766,13 @@ extern "C" __device__ float test_scalbnf(float x, int y) {
|
||||
|
||||
// DEFAULT-LABEL: @test_scalbn(
|
||||
// DEFAULT-NEXT: entry:
|
||||
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
|
||||
// DEFAULT-NEXT: ret double [[CALL_I]]
|
||||
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
|
||||
// DEFAULT-NEXT: ret double [[TMP0]]
|
||||
//
|
||||
// FINITEONLY-LABEL: @test_scalbn(
|
||||
// FINITEONLY-NEXT: entry:
|
||||
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_scalbn_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
|
||||
// FINITEONLY-NEXT: ret double [[CALL_I]]
|
||||
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
|
||||
// FINITEONLY-NEXT: ret double [[TMP0]]
|
||||
//
|
||||
extern "C" __device__ double test_scalbn(double x, int y) {
|
||||
return scalbn(x, y);
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user