Author: Matt Arsenault Date: 2023-06-23T09:11:06-04:00 New Revision: d7feba74b649cb8aabc0cb09b9bef5d79659e31c
URL: https://github.com/llvm/llvm-project/commit/d7feba74b649cb8aabc0cb09b9bef5d79659e31c DIFF: https://github.com/llvm/llvm-project/commit/d7feba74b649cb8aabc0cb09b9bef5d79659e31c.diff LOG: HIP: Directly call trunc builtins Added: Modified: clang/lib/Headers/__clang_hip_math.h clang/test/Headers/__clang_hip_math.hip Removed: ################################################################################ diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h index 850dbbef0f4e3..7e95b66232122 100644 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -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); } @@ -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); } diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index 80501d1d43e46..984adf6da4ba2 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -3073,13 +3073,13 @@ 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); @@ -3087,13 +3087,13 @@ extern "C" __device__ float test_truncf(float 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); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
