arsenm updated this revision to Diff 477238.
arsenm added a comment.
Catch another hidden in another header
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D138394/new/
https://reviews.llvm.org/D138394
Files:
clang/lib/Headers/__clang_hip_cmath.h
clang/lib/Headers/__clang_hip_math.h
clang/test/Headers/__clang_hip_cmath.hip
clang/test/Headers/__clang_hip_math.hip
Index: clang/test/Headers/__clang_hip_math.hip
===================================================================
--- clang/test/Headers/__clang_hip_math.hip
+++ clang/test/Headers/__clang_hip_math.hip
@@ -943,13 +943,13 @@
// DEFAULT-LABEL: @test_fmaf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
-// DEFAULT-NEXT: ret float [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_fmaf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
-// FINITEONLY-NEXT: ret float [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_fmaf(float x, float y, float z) {
return fmaf(x, y, z);
@@ -957,13 +957,13 @@
// DEFAULT-LABEL: @test_fma(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_fma(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_fma(double x, double y, double z) {
return fma(x, y, z);
@@ -971,13 +971,13 @@
// DEFAULT-LABEL: @test_fma_rn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_fma_rn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_fma_rn(double x, double y, double z) {
return __fma_rn(x, y, z);
@@ -3360,13 +3360,13 @@
// DEFAULT-LABEL: @test__fmaf_rn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
-// DEFAULT-NEXT: ret float [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test__fmaf_rn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
-// FINITEONLY-NEXT: ret float [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
return __fmaf_rn(x, y, z);
@@ -3624,13 +3624,13 @@
// DEFAULT-LABEL: @test__fma_rn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test__fma_rn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test__fma_rn(double x, double y, double z) {
return __fma_rn(x, y, z);
Index: clang/test/Headers/__clang_hip_cmath.hip
===================================================================
--- clang/test/Headers/__clang_hip_cmath.hip
+++ clang/test/Headers/__clang_hip_cmath.hip
@@ -18,13 +18,13 @@
// DEFAULT-LABEL: @test_fma_f16(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract half @__ocml_fma_f16(half noundef [[X:%.*]], half noundef [[Y:%.*]], half noundef [[Z:%.*]]) #[[ATTR6:[0-9]+]]
-// DEFAULT-NEXT: ret half [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
+// DEFAULT-NEXT: ret half [[TMP0]]
//
// FINITEONLY-LABEL: @test_fma_f16(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract half @__ocml_fma_f16(half noundef [[X:%.*]], half noundef [[Y:%.*]], half noundef [[Z:%.*]]) #[[ATTR6:[0-9]+]]
-// FINITEONLY-NEXT: ret half [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
+// FINITEONLY-NEXT: ret half [[TMP0]]
//
extern "C" __device__ _Float16 test_fma_f16(_Float16 x, _Float16 y,
_Float16 z) {
@@ -33,12 +33,12 @@
// DEFAULT-LABEL: @test_pow_f16(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR7:[0-9]+]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR6:[0-9]+]]
// DEFAULT-NEXT: ret half [[CALL_I]]
//
// FINITEONLY-LABEL: @test_pow_f16(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR7:[0-9]+]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR6:[0-9]+]]
// FINITEONLY-NEXT: ret half [[CALL_I]]
//
extern "C" __device__ _Float16 test_pow_f16(_Float16 x, int y) {
Index: clang/lib/Headers/__clang_hip_math.h
===================================================================
--- clang/lib/Headers/__clang_hip_math.h
+++ clang/lib/Headers/__clang_hip_math.h
@@ -255,7 +255,7 @@
__DEVICE__
float fmaf(float __x, float __y, float __z) {
- return __ocml_fma_f32(__x, __y, __z);
+ return __builtin_fmaf(__x, __y, __z);
}
__DEVICE__
@@ -633,7 +633,7 @@
#else
__DEVICE__
float __fmaf_rn(float __x, float __y, float __z) {
- return __ocml_fma_f32(__x, __y, __z);
+ return __builtin_fmaf(__x, __y, __z);
}
#endif
@@ -811,7 +811,7 @@
__DEVICE__
double fma(double __x, double __y, double __z) {
- return __ocml_fma_f64(__x, __y, __z);
+ return __builtin_fma(__x, __y, __z);
}
__DEVICE__
@@ -1270,7 +1270,7 @@
#else
__DEVICE__
double __fma_rn(double __x, double __y, double __z) {
- return __ocml_fma_f64(__x, __y, __z);
+ return __builtin_fma(__x, __y, __z);
}
#endif
// END INTRINSICS
Index: clang/lib/Headers/__clang_hip_cmath.h
===================================================================
--- clang/lib/Headers/__clang_hip_cmath.h
+++ clang/lib/Headers/__clang_hip_cmath.h
@@ -171,7 +171,7 @@
// Other functions.
__DEVICE__ __CONSTEXPR__ _Float16 fma(_Float16 __x, _Float16 __y,
_Float16 __z) {
- return __ocml_fma_f16(__x, __y, __z);
+ return __builtin_fmaf16(__x, __y, __z);
}
__DEVICE__ __CONSTEXPR__ _Float16 pow(_Float16 __base, int __iexp) {
return __ocml_pown_f16(__base, __iexp);
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
