[clang] 71be91e - HIP: Directly call rint builtins
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Tue Jul 25 04:54:16 PDT 2023
Author: Matt Arsenault
Date: 2023-07-25T07:54:11-04:00
New Revision: 71be91eba96d80d15689e4f516141c533c3c086d
URL: https://github.com/llvm/llvm-project/commit/71be91eba96d80d15689e4f516141c533c3c086d
DIFF: https://github.com/llvm/llvm-project/commit/71be91eba96d80d15689e4f516141c533c3c086d.diff
LOG: HIP: Directly call rint 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 86013d0ffc2b7a..c12777f0476f1b 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -309,7 +309,7 @@ __DEVICE__
float lgammaf(float __x) { return __ocml_lgamma_f32(__x); }
__DEVICE__
-long long int llrintf(float __x) { return __ocml_rint_f32(__x); }
+long long int llrintf(float __x) { return __builtin_rintf(__x); }
__DEVICE__
long long int llroundf(float __x) { return __builtin_roundf(__x); }
@@ -330,7 +330,7 @@ __DEVICE__
float logf(float __x) { return __builtin_logf(__x); }
__DEVICE__
-long int lrintf(float __x) { return __ocml_rint_f32(__x); }
+long int lrintf(float __x) { return __builtin_rintf(__x); }
__DEVICE__
long int lroundf(float __x) { return __builtin_roundf(__x); }
@@ -435,7 +435,7 @@ __DEVICE__
float rhypotf(float __x, float __y) { return __ocml_rhypot_f32(__x, __y); }
__DEVICE__
-float rintf(float __x) { return __ocml_rint_f32(__x); }
+float rintf(float __x) { return __builtin_rintf(__x); }
__DEVICE__
float rnorm3df(float __x, float __y, float __z) {
@@ -857,7 +857,7 @@ __DEVICE__
double lgamma(double __x) { return __ocml_lgamma_f64(__x); }
__DEVICE__
-long long int llrint(double __x) { return __ocml_rint_f64(__x); }
+long long int llrint(double __x) { return __builtin_rint(__x); }
__DEVICE__
long long int llround(double __x) { return __builtin_round(__x); }
@@ -878,7 +878,7 @@ __DEVICE__
double logb(double __x) { return __ocml_logb_f64(__x); }
__DEVICE__
-long int lrint(double __x) { return __ocml_rint_f64(__x); }
+long int lrint(double __x) { return __builtin_rint(__x); }
__DEVICE__
long int lround(double __x) { return __builtin_round(__x); }
@@ -991,7 +991,7 @@ __DEVICE__
double rhypot(double __x, double __y) { return __ocml_rhypot_f64(__x, __y); }
__DEVICE__
-double rint(double __x) { return __ocml_rint_f64(__x); }
+double rint(double __x) { return __builtin_rint(__x); }
__DEVICE__
double rnorm(int __dim,
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index bfb66ce27e6171..aa31654deccc43 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -1479,14 +1479,14 @@ extern "C" __device__ double test_lgamma(double x) {
// DEFAULT-LABEL: @test_llrintf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
+// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_llrintf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float [[X:%.*]])
+// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
extern "C" __device__ long long int test_llrintf(float x) {
@@ -1495,14 +1495,14 @@ extern "C" __device__ long long int test_llrintf(float x) {
// DEFAULT-LABEL: @test_llrint(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
+// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_llrint(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double [[X:%.*]])
+// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
extern "C" __device__ long long int test_llrint(double x) {
@@ -1655,14 +1655,14 @@ extern "C" __device__ double test_logb(double x) {
// DEFAULT-LABEL: @test_lrintf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
+// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_lrintf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float [[X:%.*]])
+// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
extern "C" __device__ long int test_lrintf(float x) {
@@ -1671,14 +1671,14 @@ extern "C" __device__ long int test_lrintf(float x) {
// DEFAULT-LABEL: @test_lrint(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
+// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_lrint(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double [[X:%.*]])
+// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
extern "C" __device__ long int test_lrint(double x) {
@@ -2440,13 +2440,13 @@ extern "C" __device__ double test_rhypot(double x, double y) {
// DEFAULT-LABEL: @test_rintf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret float [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_rintf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret float [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float [[X:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_rintf(float x) {
return rintf(x);
@@ -2454,13 +2454,13 @@ extern "C" __device__ float test_rintf(float x) {
// DEFAULT-LABEL: @test_rint(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_rint(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double [[X:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_rint(double x) {
return rint(x);
More information about the cfe-commits
mailing list