Author: Matt Arsenault Date: 2023-07-07T09:26:07-04:00 New Revision: 9a97e3bf8c3f576b29dfcd76925b6c23b3961b99
URL: https://github.com/llvm/llvm-project/commit/9a97e3bf8c3f576b29dfcd76925b6c23b3961b99 DIFF: https://github.com/llvm/llvm-project/commit/9a97e3bf8c3f576b29dfcd76925b6c23b3961b99.diff LOG: HIP: Directly call ceil 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 a914496cb7b14a..c67f20873253b0 100644 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -182,7 +182,7 @@ __DEVICE__ float cbrtf(float __x) { return __ocml_cbrt_f32(__x); } __DEVICE__ -float ceilf(float __x) { return __ocml_ceil_f32(__x); } +float ceilf(float __x) { return __builtin_ceilf(__x); } __DEVICE__ float copysignf(float __x, float __y) { return __builtin_copysignf(__x, __y); } @@ -731,7 +731,7 @@ __DEVICE__ double cbrt(double __x) { return __ocml_cbrt_f64(__x); } __DEVICE__ -double ceil(double __x) { return __ocml_ceil_f64(__x); } +double ceil(double __x) { return __builtin_ceil(__x); } __DEVICE__ double copysign(double __x, double __y) { diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index 8e5201a89ceca2..a6797b01f25b27 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -475,13 +475,13 @@ extern "C" __device__ double test_cbrt(double x) { // DEFAULT-LABEL: @test_ceilf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR14]] -// DEFAULT-NEXT: ret float [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.ceil.f32(float [[X:%.*]]) +// DEFAULT-NEXT: ret float [[TMP0]] // // FINITEONLY-LABEL: @test_ceilf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ceil_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] -// FINITEONLY-NEXT: ret float [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.ceil.f32(float [[X:%.*]]) +// FINITEONLY-NEXT: ret float [[TMP0]] // extern "C" __device__ float test_ceilf(float x) { return ceilf(x); @@ -489,13 +489,13 @@ extern "C" __device__ float test_ceilf(float x) { // DEFAULT-LABEL: @test_ceil( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR14]] -// DEFAULT-NEXT: ret double [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.ceil.f64(double [[X:%.*]]) +// DEFAULT-NEXT: ret double [[TMP0]] // // FINITEONLY-LABEL: @test_ceil( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ceil_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] -// FINITEONLY-NEXT: ret double [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.ceil.f64(double [[X:%.*]]) +// FINITEONLY-NEXT: ret double [[TMP0]] // extern "C" __device__ double test_ceil(double x) { return ceil(x); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
