Skip to content

Commit

Permalink
HIP: Directly call trunc builtins
Browse files Browse the repository at this point in the history
  • Loading branch information
arsenm committed Jun 23, 2023
1 parent 2449931 commit d7feba7
Show file tree
Hide file tree
Showing 2 changed files with 10 additions and 10 deletions.
4 changes: 2 additions & 2 deletions clang/lib/Headers/__clang_hip_math.h
Original file line number Diff line number Diff line change
Expand Up @@ -522,7 +522,7 @@ __DEVICE__
float tgammaf(float __x) { return __ocml_tgamma_f32(__x); }

__DEVICE__
float truncf(float __x) { return __ocml_trunc_f32(__x); }
float truncf(float __x) { return __builtin_truncf(__x); }

__DEVICE__
float y0f(float __x) { return __ocml_y0_f32(__x); }
Expand Down Expand Up @@ -1078,7 +1078,7 @@ __DEVICE__
double tgamma(double __x) { return __ocml_tgamma_f64(__x); }

__DEVICE__
double trunc(double __x) { return __ocml_trunc_f64(__x); }
double trunc(double __x) { return __builtin_trunc(__x); }

__DEVICE__
double y0(double __x) { return __ocml_y0_f64(__x); }
Expand Down
16 changes: 8 additions & 8 deletions clang/test/Headers/__clang_hip_math.hip
Original file line number Diff line number Diff line change
Expand Up @@ -3073,27 +3073,27 @@ extern "C" __device__ double test_tgamma(double x) {

// DEFAULT-LABEL: @test_truncf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: ret float [[CALL_I]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.trunc.f32(float [[X:%.*]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_truncf(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_trunc_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.trunc.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_truncf(float x) {
return truncf(x);
}

// DEFAULT-LABEL: @test_trunc(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: ret double [[CALL_I]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.trunc.f64(double [[X:%.*]])
// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_trunc(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_trunc_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.trunc.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_trunc(double x) {
return trunc(x);
Expand Down

0 comments on commit d7feba7

Please sign in to comment.