Skip to content

Commit

Permalink
HIP: Directly call ldexp builtins
Browse files Browse the repository at this point in the history
  • Loading branch information
@arsenm
arsenm committed Jun 22, 2023
1 parent f14c8eb commit c448e1d
Show file tree
Hide file tree
Showing 2 changed files with 34 additions and 34 deletions.
12 changes: 6 additions & 6 deletions clang/lib/Headers/__clang_hip_math.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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); }
Expand Down Expand Up @@ -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); }
Expand Down Expand Up @@ -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); }
Expand Down Expand Up @@ -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); }
Expand Down
56 changes: 28 additions & 28 deletions clang/test/Headers/__clang_hip_math.hip
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1439,27 +1439,27 @@ 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);
}

// 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);
Expand Down Expand Up @@ -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(
Expand All @@ -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) {
Expand All @@ -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(
Expand All @@ -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) {
Expand All @@ -2752,27 +2752,27 @@ 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);
}

// 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);
Expand Down

0 comments on commit c448e1d

Please sign in to comment.