aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorMatt Arsenault <Matthew.Arsenault@amd.com>2022-11-21 23:37:15 -0500
committerMatt Arsenault <Matthew.Arsenault@amd.com>2023-07-25 07:56:31 -0400
commite7ab6982de87b14c9584e1267cd755561b4c063c (patch)
tree2d2555855ac577cc3e64f6453b5eb58cefaad3a7
parent395cd33ba850989209834a2e332d21b42168cfaf (diff)
downloadllvm-e7ab6982de87b14c9584e1267cd755561b4c063c.zip
llvm-e7ab6982de87b14c9584e1267cd755561b4c063c.tar.gz
llvm-e7ab6982de87b14c9584e1267cd755561b4c063c.tar.bz2
HIP: Directly call nearbyint builtins
-rw-r--r--clang/lib/Headers/__clang_hip_math.h4
-rw-r--r--clang/test/Headers/__clang_hip_math.hip16
2 files changed, 10 insertions, 10 deletions
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index c12777f..a47dda3 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -369,7 +369,7 @@ float nanf(const char *__tagp __attribute__((nonnull))) {
}
__DEVICE__
-float nearbyintf(float __x) { return __ocml_nearbyint_f32(__x); }
+float nearbyintf(float __x) { return __builtin_nearbyintf(__x); }
__DEVICE__
float nextafterf(float __x, float __y) {
@@ -925,7 +925,7 @@ double nan(const char *__tagp) {
}
__DEVICE__
-double nearbyint(double __x) { return __ocml_nearbyint_f64(__x); }
+double nearbyint(double __x) { return __builtin_nearbyint(__x); }
__DEVICE__
double nextafter(double __x, double __y) {
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index aa31654..85d559f 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -1996,13 +1996,13 @@ extern "C" __device__ double test_nan_fill() {
// DEFAULT-LABEL: @test_nearbyintf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret float [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.nearbyint.f32(float [[X:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_nearbyintf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_nearbyint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret float [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.nearbyint.f32(float [[X:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_nearbyintf(float x) {
return nearbyintf(x);
@@ -2010,13 +2010,13 @@ extern "C" __device__ float test_nearbyintf(float x) {
// DEFAULT-LABEL: @test_nearbyint(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.nearbyint.f64(double [[X:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_nearbyint(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_nearbyint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.nearbyint.f64(double [[X:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_nearbyint(double x) {
return nearbyint(x);