[clang] clang/HIP: Add tests that shows fpmath metadata ends up on sqrt calls (PR #136413)
via cfe-commits
cfe-commits at lists.llvm.org
Sat Apr 19 00:47:09 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Matt Arsenault (arsenm)
<details>
<summary>Changes</summary>
Make sure the builtin header sqrts work with
-fno-hip-f32-correctly-rounded-divide-sqrt, and we end up with
properly annotated sqrt intrinsic callsites.
---
Patch is 169.18 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/136413.diff
2 Files Affected:
- (modified) clang/test/Headers/__clang_hip_cmath.hip (+43-1)
- (modified) clang/test/Headers/__clang_hip_math.hip (+1439)
``````````diff
diff --git a/clang/test/Headers/__clang_hip_cmath.hip b/clang/test/Headers/__clang_hip_cmath.hip
index 7d812fd0265a6..fcd74996e5fa4 100644
--- a/clang/test/Headers/__clang_hip_cmath.hip
+++ b/clang/test/Headers/__clang_hip_cmath.hip
@@ -6,7 +6,7 @@
// RUN: -internal-isystem %S/Inputs/include \
// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \
-// RUN: -D__HIPCC_RTC__ | FileCheck -check-prefix=DEFAULT %s
+// RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=DEFAULT,CORRECT-DIV-SQRT %s
// Check that we end up with fast math flags set on intrinsic calls
// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
@@ -17,6 +17,15 @@
// RUN: -menable-no-nans -o - \
// RUN: -D__HIPCC_RTC__ | FileCheck -check-prefix=FINITEONLY %s
+// Check that we end up with fpmath metadata set on sqrt calls
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN: -internal-isystem %S/Inputs/include \
+// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 \
+// RUN: -fno-hip-fp32-correctly-rounded-divide-sqrt -o - \
+// RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=DEFAULT,NO-CORRECT-DIV-SQRT %s
+
// DEFAULT-LABEL: @test_fma_f16(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
@@ -142,3 +151,36 @@ namespace user_namespace {
fma(a, b, b);
}
}
+
+// CORRECT-DIV-SQRT-LABEL: @test_sqrt_f32(
+// CORRECT-DIV-SQRT-NEXT: entry:
+// CORRECT-DIV-SQRT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X:%.*]])
+// CORRECT-DIV-SQRT-NEXT: ret float [[TMP0]]
+//
+// FINITEONLY-LABEL: @test_sqrt_f32(
+// FINITEONLY-NEXT: entry:
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.sqrt.f32(float nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
+//
+// NO-CORRECT-DIV-SQRT-LABEL: @test_sqrt_f32(
+// NO-CORRECT-DIV-SQRT-NEXT: entry:
+// NO-CORRECT-DIV-SQRT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X:%.*]]), !fpmath [[META4:![0-9]+]]
+// NO-CORRECT-DIV-SQRT-NEXT: ret float [[TMP0]]
+//
+extern "C" __device__ float test_sqrt_f32(float x) {
+ return sqrt(x);
+}
+
+// DEFAULT-LABEL: @test_sqrt_f64(
+// DEFAULT-NEXT: entry:
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[X:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
+//
+// FINITEONLY-LABEL: @test_sqrt_f64(
+// FINITEONLY-NEXT: entry:
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
+//
+extern "C" __device__ double test_sqrt_f64(double x) {
+ return sqrt(x);
+}
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index df1cd716342a5..11c9cd301abb7 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -27,6 +27,15 @@
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -fgpu-approx-transcendentals -o - \
// RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,APPROX %s
+// Check that we end up with fpmath metadata set on sqrt calls
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN: -internal-isystem %S/Inputs/include \
+// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -fno-hip-fp32-correctly-rounded-divide-sqrt -o - \
+// RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,NCRDIV %s
+
+
// Check that we use the AMDGCNSPIRV address space map
// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
@@ -465,6 +474,11 @@ extern "C" __device__ long long test_llabs(long x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR12:[0-9]+]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_acosf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR12:[0-9]+]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_acosf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR12:[0-9]+]]
@@ -489,6 +503,11 @@ extern "C" __device__ float test_acosf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_acos(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_acos(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR12]]
@@ -513,6 +532,11 @@ extern "C" __device__ double test_acos(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR13:[0-9]+]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_acoshf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR13:[0-9]+]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_acoshf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR13:[0-9]+]]
@@ -537,6 +561,11 @@ extern "C" __device__ float test_acoshf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_acosh(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_acosh(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -561,6 +590,11 @@ extern "C" __device__ double test_acosh(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_asinf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_asinf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR12]]
@@ -585,6 +619,11 @@ extern "C" __device__ float test_asinf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_asin(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_asin(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR12]]
@@ -610,6 +649,11 @@ extern "C" __device__ double test_asin(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_asinhf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_asinhf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -634,6 +678,11 @@ extern "C" __device__ float test_asinhf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_asinh(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_asinh(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -658,6 +707,11 @@ extern "C" __device__ double test_asinh(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_atan2f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_atan2f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
@@ -682,6 +736,11 @@ extern "C" __device__ float test_atan2f(float x, float y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_atan2(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_atan2(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
@@ -706,6 +765,11 @@ extern "C" __device__ double test_atan2(double x, double y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_atanf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_atanf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR12]]
@@ -730,6 +794,11 @@ extern "C" __device__ float test_atanf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_atan(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_atan(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR12]]
@@ -754,6 +823,11 @@ extern "C" __device__ double test_atan(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_atanhf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_atanhf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -778,6 +852,11 @@ extern "C" __device__ float test_atanhf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_atanh(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_atanh(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -802,6 +881,11 @@ extern "C" __device__ double test_atanh(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_cbrtf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_cbrtf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -826,6 +910,11 @@ extern "C" __device__ float test_cbrtf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_cbrt(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_cbrt(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -850,6 +939,11 @@ extern "C" __device__ double test_cbrt(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.ceil.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_ceilf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.ceil.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_ceilf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.ceil.f32(float [[X:%.*]])
@@ -874,6 +968,11 @@ extern "C" __device__ float test_ceilf(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.ceil.f64(double [[X:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_ceil(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.ceil.f64(double [[X:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_ceil(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.ceil.f64(double [[X:%.*]])
@@ -898,6 +997,11 @@ extern "C" __device__ double test_ceil(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.copysign.f32(float [[X:%.*]], float [[Y:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_copysignf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.copysign.f32(float [[X:%.*]], float [[Y:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_copysignf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.copysign.f32(float [[X:%.*]], float [[Y:%.*]])
@@ -922,6 +1026,11 @@ extern "C" __device__ float test_copysignf(float x, float y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.copysign.f64(double [[X:%.*]], double [[Y:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_copysign(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.copysign.f64(double [[X:%.*]], double [[Y:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_copysign(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.copysign.f64(double [[X:%.*]], double [[Y:%.*]])
@@ -946,6 +1055,11 @@ extern "C" __device__ double test_copysign(double x, double y) {
// APPROX-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
// APPROX-NEXT: ret float [[CALL_I1]]
//
+// NCRDIV-LABEL: @test_cosf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_cosf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
@@ -970,6 +1084,11 @@ extern "C" __device__ float test_cosf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_cos(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_cos(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -994,6 +1113,11 @@ extern "C" __device__ double test_cos(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_coshf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_coshf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -1018,6 +1142,11 @@ extern "C" __device__ float test_coshf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_cosh(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEX...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/136413
More information about the cfe-commits
mailing list