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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
