Author: Matt Arsenault Date: 2023-06-28T16:57:53-04:00 New Revision: fe65043a67c85345887339f371cca1a0c735a639
URL: https://github.com/llvm/llvm-project/commit/fe65043a67c85345887339f371cca1a0c735a639 DIFF: https://github.com/llvm/llvm-project/commit/fe65043a67c85345887339f371cca1a0c735a639.diff LOG: HIP: Directly call floor builtins Added: Modified: clang/lib/Headers/__clang_hip_math.h clang/test/Headers/__clang_hip_math.hip clang/test/Headers/hip-header.hip Removed: ################################################################################ diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h index 26c2f77c82c63..a914496cb7b14 100644 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -239,7 +239,7 @@ __DEVICE__ float fdividef(float __x, float __y) { return __x / __y; } __DEVICE__ -float floorf(float __x) { return __ocml_floor_f32(__x); } +float floorf(float __x) { return __builtin_floorf(__x); } __DEVICE__ float fmaf(float __x, float __y, float __z) { @@ -787,7 +787,7 @@ __DEVICE__ double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); } __DEVICE__ -double floor(double __x) { return __ocml_floor_f64(__x); } +double floor(double __x) { return __builtin_floor(__x); } __DEVICE__ double fma(double __x, double __y, double __z) { diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index e507aa3ccf4ef..8e5201a89ceca 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -909,13 +909,13 @@ extern "C" __device__ float test_fdividef(float x, float y) { // DEFAULT-LABEL: @test_floorf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR14]] -// DEFAULT-NEXT: ret float [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.floor.f32(float [[X:%.*]]) +// DEFAULT-NEXT: ret float [[TMP0]] // // FINITEONLY-LABEL: @test_floorf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_floor_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] -// FINITEONLY-NEXT: ret float [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.floor.f32(float [[X:%.*]]) +// FINITEONLY-NEXT: ret float [[TMP0]] // extern "C" __device__ float test_floorf(float x) { return floorf(x); @@ -923,13 +923,13 @@ extern "C" __device__ float test_floorf(float x) { // DEFAULT-LABEL: @test_floor( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR14]] -// DEFAULT-NEXT: ret double [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.floor.f64(double [[X:%.*]]) +// DEFAULT-NEXT: ret double [[TMP0]] // // FINITEONLY-LABEL: @test_floor( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_floor_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] -// FINITEONLY-NEXT: ret double [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.floor.f64(double [[X:%.*]]) +// FINITEONLY-NEXT: ret double [[TMP0]] // extern "C" __device__ double test_floor(double x) { return floor(x); diff --git a/clang/test/Headers/hip-header.hip b/clang/test/Headers/hip-header.hip index 38567cb3b202f..7d7b5e7176d9d 100644 --- a/clang/test/Headers/hip-header.hip +++ b/clang/test/Headers/hip-header.hip @@ -103,7 +103,7 @@ __device__ size_t test_size_t() { // Check there is no ambiguity when calling overloaded math functions. // CHECK-LABEL: define{{.*}}@_Z10test_floorv -// CHECK: call {{.*}}double @__ocml_floor_f64(double +// CHECK: call {{.*}}double @llvm.floor.f64(double __device__ float test_floor() { return floor(5); } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
