[clang] 9a31751 - HIP: Directly call round builtins
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 7 06:29:36 PDT 2023
Author: Matt Arsenault
Date: 2023-07-07T09:26:40-04:00
New Revision: 9a317516a515f3c5b15f9060329a503e8f261c7f
URL: https://github.com/llvm/llvm-project/commit/9a317516a515f3c5b15f9060329a503e8f261c7f
DIFF: https://github.com/llvm/llvm-project/commit/9a317516a515f3c5b15f9060329a503e8f261c7f.diff
LOG: HIP: Directly call round 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 c67f20873253b0..8a90c4acd94105 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -312,7 +312,7 @@ __DEVICE__
long long int llrintf(float __x) { return __ocml_rint_f32(__x); }
__DEVICE__
-long long int llroundf(float __x) { return __ocml_round_f32(__x); }
+long long int llroundf(float __x) { return __builtin_roundf(__x); }
__DEVICE__
float log10f(float __x) { return __ocml_log10_f32(__x); }
@@ -333,7 +333,7 @@ __DEVICE__
long int lrintf(float __x) { return __ocml_rint_f32(__x); }
__DEVICE__
-long int lroundf(float __x) { return __ocml_round_f32(__x); }
+long int lroundf(float __x) { return __builtin_roundf(__x); }
__DEVICE__
float modff(float __x, float *__iptr) {
@@ -460,7 +460,7 @@ float rnormf(int __dim,
}
__DEVICE__
-float roundf(float __x) { return __ocml_round_f32(__x); }
+float roundf(float __x) { return __builtin_roundf(__x); }
__DEVICE__
float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); }
@@ -860,7 +860,7 @@ __DEVICE__
long long int llrint(double __x) { return __ocml_rint_f64(__x); }
__DEVICE__
-long long int llround(double __x) { return __ocml_round_f64(__x); }
+long long int llround(double __x) { return __builtin_round(__x); }
__DEVICE__
double log(double __x) { return __ocml_log_f64(__x); }
@@ -881,7 +881,7 @@ __DEVICE__
long int lrint(double __x) { return __ocml_rint_f64(__x); }
__DEVICE__
-long int lround(double __x) { return __ocml_round_f64(__x); }
+long int lround(double __x) { return __builtin_round(__x); }
__DEVICE__
double modf(double __x, double *__iptr) {
@@ -1016,7 +1016,7 @@ double rnorm4d(double __x, double __y, double __z, double __w) {
}
__DEVICE__
-double round(double __x) { return __ocml_round_f64(__x); }
+double round(double __x) { return __builtin_round(__x); }
__DEVICE__
double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); }
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index a6797b01f25b27..35f637b912f439 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -1515,14 +1515,14 @@ extern "C" __device__ long long int test_llrint(double x) {
// DEFAULT-LABEL: @test_llroundf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]])
+// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_llroundf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_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.round.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_llroundf(float x) {
@@ -1531,14 +1531,14 @@ extern "C" __device__ long long int test_llroundf(float x) {
// DEFAULT-LABEL: @test_llround(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]])
+// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_llround(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_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.round.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_llround(double x) {
@@ -1691,14 +1691,14 @@ extern "C" __device__ long int test_lrint(double x) {
// DEFAULT-LABEL: @test_lroundf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]])
+// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_lroundf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_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.round.f32(float [[X:%.*]])
+// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
extern "C" __device__ long int test_lroundf(float x) {
@@ -1707,14 +1707,14 @@ extern "C" __device__ long int test_lroundf(float x) {
// DEFAULT-LABEL: @test_lround(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]])
+// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_lround(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_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.round.f64(double [[X:%.*]])
+// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
extern "C" __device__ long int test_lround(double x) {
@@ -2616,13 +2616,13 @@ extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w
// DEFAULT-LABEL: @test_roundf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret float [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_roundf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret float [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float [[X:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_roundf(float x) {
return roundf(x);
@@ -2630,13 +2630,13 @@ extern "C" __device__ float test_roundf(float x) {
// DEFAULT-LABEL: @test_round(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_round(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double [[X:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_round(double x) {
return round(x);
More information about the cfe-commits
mailing list