[clang] 4524a31 - HIP: Directly call copysign builtins
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Thu Jun 8 13:16:16 PDT 2023
Author: Matt Arsenault
Date: 2023-06-08T16:16:11-04:00
New Revision: 4524a3159e8407d2b453293eff694f3a1534462d
URL: https://github.com/llvm/llvm-project/commit/4524a3159e8407d2b453293eff694f3a1534462d
DIFF: https://github.com/llvm/llvm-project/commit/4524a3159e8407d2b453293eff694f3a1534462d.diff
LOG: HIP: Directly call copysign 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 40ee1598200b5..c19e32bd29364 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -185,7 +185,7 @@ __DEVICE__
float ceilf(float __x) { return __ocml_ceil_f32(__x); }
__DEVICE__
-float copysignf(float __x, float __y) { return __ocml_copysign_f32(__x, __y); }
+float copysignf(float __x, float __y) { return __builtin_copysignf(__x, __y); }
__DEVICE__
float cosf(float __x) { return __ocml_cos_f32(__x); }
@@ -736,7 +736,7 @@ double ceil(double __x) { return __ocml_ceil_f64(__x); }
__DEVICE__
double copysign(double __x, double __y) {
- return __ocml_copysign_f64(__x, __y);
+ return __builtin_copysign(__x, __y);
}
__DEVICE__
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index 453acda646a95..1b4cfa3ab31e7 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -503,13 +503,13 @@ extern "C" __device__ double test_ceil(double x) {
// DEFAULT-LABEL: @test_copysignf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret float [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.copysign.f32(float [[X:%.*]], float [[Y:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_copysignf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_copysign_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret float [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.copysign.f32(float [[X:%.*]], float [[Y:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_copysignf(float x, float y) {
return copysignf(x, y);
@@ -517,13 +517,13 @@ extern "C" __device__ float test_copysignf(float x, float y) {
// DEFAULT-LABEL: @test_copysign(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.copysign.f64(double [[X:%.*]], double [[Y:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_copysign(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_copysign_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.copysign.f64(double [[X:%.*]], double [[Y:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_copysign(double x, double y) {
return copysign(x, y);
More information about the cfe-commits
mailing list