Skip to content

clang/HIP: Add tests that shows fpmath metadata ends up on sqrt calls #136413

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Merged
merged 1 commit into from
Apr 24, 2025
Merged

clang/HIP: Add tests that shows fpmath metadata ends up on sqrt calls #136413

merged 1 commit into from
Apr 24, 2025

Conversation

arsenm
Copy link
Contributor

@arsenm arsenm commented Apr 19, 2025

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.

@arsenm
clang/HIP: Add tests that shows fpmath metadata ends up on sqrt calls
577796c 
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.
@arsenm Graphite App
Copy link
Contributor Author

arsenm commented Apr 19, 2025

This stack of pull requests is managed by Graphite. Learn more about stacking.

@arsenm arsenm requested a review from yxsamliu April 19, 2025 07:46
@arsenm arsenm marked this pull request as ready for review April 19, 2025 07:46
@llvmbot llvmbot added the clang Clang issues not falling into any other category label Apr 19, 2025
@llvmbot
Copy link
Member

llvmbot commented Apr 19, 2025 •
edited
Loading

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-backend-amdgpu

Author: Matt Arsenault (arsenm)

Changes

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 --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]

@arsenm
Copy link
Contributor Author

arsenm commented Apr 24, 2025

ping

yxsamliu
yxsamliu reviewed Apr 24, 2025
View reviewed changes
Copy link
Collaborator

@yxsamliu yxsamliu left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM. Thanks

@arsenm arsenm merged commit 3e7e23d into main Apr 24, 2025
16 checks passed
@arsenm arsenm deleted the users/arsenm/clang/add-hip-tests-for-sqrt-metadata branch April 24, 2025 17:48
IanWood1 pushed a commit to IanWood1/llvm-project that referenced this pull request May 6, 2025
@arsenm @IanWood1
clang/HIP: Add tests that shows fpmath metadata ends up on sqrt calls (…
1da518f 
…llvm#136413)

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.
IanWood1 pushed a commit to IanWood1/llvm-project that referenced this pull request May 6, 2025
@arsenm @IanWood1
clang/HIP: Add tests that shows fpmath metadata ends up on sqrt calls (…
07d9636 
…llvm#136413)

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.
IanWood1 pushed a commit to IanWood1/llvm-project that referenced this pull request May 6, 2025
@arsenm @IanWood1
clang/HIP: Add tests that shows fpmath metadata ends up on sqrt calls (…
03ccc01 
…llvm#136413)

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.
Ankur-0429 pushed a commit to Ankur-0429/llvm-project that referenced this pull request May 9, 2025
@arsenm @Ankur-0429
clang/HIP: Add tests that shows fpmath metadata ends up on sqrt calls (…
1061f23 
…llvm#136413)

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.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@yxsamliu yxsamliu yxsamliu approved these changes

@grantchen08 grantchen08 grantchen08 approved these changes

Assignees
No one assigned
Labels
backend:AMDGPU clang Clang issues not falling into any other category
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
4 participants
@arsenm @llvmbot @yxsamliu @grantchen08