[clang] clang/HIP: Add tests that shows fpmath metadata ends up on sqrt calls (PR #136413)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Sat Apr 19 00:46:19 PDT 2025
https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/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.
>From 577796c23b8e65dcdd2b45767386a81ff70f4e84 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Sat, 19 Apr 2025 09:44:26 +0200
Subject: [PATCH] clang/HIP: Add tests that shows fpmath metadata ends up on
sqrt calls
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.
---
clang/test/Headers/__clang_hip_cmath.hip | 44 +-
clang/test/Headers/__clang_hip_math.hip | 1439 ++++++++++++++++++++++
2 files changed, 1482 insertions(+), 1 deletion(-)
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-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_cosh(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -1042,6 +1171,11 @@ extern "C" __device__ double test_cosh(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_cospif(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_cospif(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -1066,6 +1200,11 @@ extern "C" __device__ float test_cospif(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_cospi(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_cospi(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -1090,6 +1229,11 @@ extern "C" __device__ double test_cospi(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_cyl_bessel_i0f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_cyl_bessel_i0f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -1114,6 +1258,11 @@ extern "C" __device__ float test_cyl_bessel_i0f(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_cyl_bessel_i0(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_cyl_bessel_i0(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -1138,6 +1287,11 @@ extern "C" __device__ double test_cyl_bessel_i0(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_cyl_bessel_i1f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_cyl_bessel_i1f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -1162,6 +1316,11 @@ extern "C" __device__ float test_cyl_bessel_i1f(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_cyl_bessel_i1(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_cyl_bessel_i1(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -1186,6 +1345,11 @@ extern "C" __device__ double test_cyl_bessel_i1(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_erfcf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_erfcf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -1210,6 +1374,11 @@ extern "C" __device__ float test_erfcf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_erfc(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_erfc(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -1234,6 +1403,11 @@ extern "C" __device__ double test_erfc(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_erfinvf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_erfinvf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -1258,6 +1432,11 @@ extern "C" __device__ float test_erfinvf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_erfinv(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_erfinv(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -1282,6 +1461,11 @@ extern "C" __device__ double test_erfinv(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.exp10.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_exp10f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.exp10.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_exp10f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.exp10.f32(float [[X:%.*]])
@@ -1306,6 +1490,11 @@ extern "C" __device__ float test_exp10f(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_exp10(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_exp10(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -1330,6 +1519,11 @@ extern "C" __device__ double test_exp10(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.exp2.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_exp2f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.exp2.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_exp2f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.exp2.f32(float [[X:%.*]])
@@ -1354,6 +1548,11 @@ extern "C" __device__ float test_exp2f(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_exp2(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_exp2(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -1378,6 +1577,11 @@ extern "C" __device__ double test_exp2(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.exp.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_expf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.exp.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_expf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.exp.f32(float [[X:%.*]])
@@ -1402,6 +1606,11 @@ extern "C" __device__ float test_expf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_exp(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_exp(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -1426,6 +1635,11 @@ extern "C" __device__ double test_exp(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_expm1f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_expm1f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -1450,6 +1664,11 @@ extern "C" __device__ float test_expm1f(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_expm1(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_expm1(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -1474,6 +1693,11 @@ extern "C" __device__ double test_expm1(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.fabs.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_fabsf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.fabs.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_fabsf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.fabs.f32(float [[X:%.*]])
@@ -1498,6 +1722,11 @@ extern "C" __device__ float test_fabsf(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.fabs.f64(double [[X:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_fabs(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.fabs.f64(double [[X:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_fabs(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.fabs.f64(double [[X:%.*]])
@@ -1522,6 +1751,11 @@ extern "C" __device__ double test_fabs(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_fdimf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_fdimf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
@@ -1546,6 +1780,11 @@ extern "C" __device__ float test_fdimf(float x, float y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_fdim(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_fdim(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
@@ -1570,6 +1809,11 @@ extern "C" __device__ double test_fdim(double x, double y) {
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]]
// APPROX-NEXT: ret float [[DIV_I]]
//
+// NCRDIV-LABEL: @test_fdividef(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]], !fpmath [[META12:![0-9]+]]
+// NCRDIV-NEXT: ret float [[DIV_I]]
+//
// AMDGCNSPIRV-LABEL: @test_fdividef(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]]
@@ -1594,6 +1838,11 @@ extern "C" __device__ float test_fdividef(float x, float y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.floor.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_floorf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.floor.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_floorf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.floor.f32(float [[X:%.*]])
@@ -1618,6 +1867,11 @@ extern "C" __device__ float test_floorf(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.floor.f64(double [[X:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_floor(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.floor.f64(double [[X:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_floor(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.floor.f64(double [[X:%.*]])
@@ -1642,6 +1896,11 @@ extern "C" __device__ double test_floor(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_fmaf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_fmaf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
@@ -1666,6 +1925,11 @@ extern "C" __device__ float test_fmaf(float x, float y, float z) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_fma(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_fma(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
@@ -1690,6 +1954,11 @@ extern "C" __device__ double test_fma(double x, double y, double z) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_fma_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_fma_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
@@ -1714,6 +1983,11 @@ extern "C" __device__ double test_fma_rn(double x, double y, double z) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_fmaxf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_fmaxf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
@@ -1738,6 +2012,11 @@ extern "C" __device__ float test_fmaxf(float x, float y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_fmax(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_fmax(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
@@ -1762,6 +2041,11 @@ extern "C" __device__ double test_fmax(double x, double y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_fminf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_fminf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
@@ -1786,6 +2070,11 @@ extern "C" __device__ float test_fminf(float x, float y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_fmin(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_fmin(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
@@ -1810,6 +2099,11 @@ extern "C" __device__ double test_fmin(double x, double y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_fmodf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_fmodf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
@@ -1834,6 +2128,11 @@ extern "C" __device__ float test_fmodf(float x, float y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_fmod(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_fmod(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
@@ -1867,6 +2166,14 @@ extern "C" __device__ double test_fmod(double x, double y) {
// APPROX-NEXT: [[TMP2:%.*]] = extractvalue { float, i32 } [[TMP0]], 0
// APPROX-NEXT: ret float [[TMP2]]
//
+// NCRDIV-LABEL: @test_frexpf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X:%.*]])
+// NCRDIV-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1
+// NCRDIV-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13:![0-9]+]]
+// NCRDIV-NEXT: [[TMP2:%.*]] = extractvalue { float, i32 } [[TMP0]], 0
+// NCRDIV-NEXT: ret float [[TMP2]]
+//
// AMDGCNSPIRV-LABEL: @test_frexpf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call addrspace(4) { float, i32 } @llvm.frexp.f32.i32(float [[X:%.*]])
@@ -1903,6 +2210,14 @@ extern "C" __device__ float test_frexpf(float x, int* y) {
// APPROX-NEXT: [[TMP2:%.*]] = extractvalue { double, i32 } [[TMP0]], 0
// APPROX-NEXT: ret double [[TMP2]]
//
+// NCRDIV-LABEL: @test_frexp(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X:%.*]])
+// NCRDIV-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1
+// NCRDIV-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
+// NCRDIV-NEXT: [[TMP2:%.*]] = extractvalue { double, i32 } [[TMP0]], 0
+// NCRDIV-NEXT: ret double [[TMP2]]
+//
// AMDGCNSPIRV-LABEL: @test_frexp(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call addrspace(4) { double, i32 } @llvm.frexp.f64.i32(double [[X:%.*]])
@@ -1930,6 +2245,11 @@ extern "C" __device__ double test_frexp(double x, int* y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_hypotf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_hypotf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
@@ -1954,6 +2274,11 @@ extern "C" __device__ float test_hypotf(float x, float y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_hypot(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_hypot(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
@@ -1978,6 +2303,11 @@ extern "C" __device__ double test_hypot(double x, double y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call noundef i32 @__ocml_ilogb_f32(float noundef [[X:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret i32 [[CALL_I]]
//
+// NCRDIV-LABEL: @test_ilogbf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call noundef i32 @__ocml_ilogb_f32(float noundef [[X:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret i32 [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_ilogbf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call spir_func noundef addrspace(4) i32 @__ocml_ilogb_f32(float noundef [[X:%.*]]) #[[ATTR12]]
@@ -2002,6 +2332,11 @@ extern "C" __device__ int test_ilogbf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call noundef i32 @__ocml_ilogb_f64(double noundef [[X:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret i32 [[CALL_I]]
//
+// NCRDIV-LABEL: @test_ilogb(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call noundef i32 @__ocml_ilogb_f64(double noundef [[X:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret i32 [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_ilogb(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call spir_func noundef addrspace(4) i32 @__ocml_ilogb_f64(double noundef [[X:%.*]]) #[[ATTR12]]
@@ -2029,6 +2364,13 @@ extern "C" __device__ int test_ilogb(double x) {
// APPROX-NEXT: [[CONV:%.*]] = zext i1 [[TMP1]] to i32
// APPROX-NEXT: ret i32 [[CONV]]
//
+// NCRDIV-LABEL: @test___finitef(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call float @llvm.fabs.f32(float [[X:%.*]])
+// NCRDIV-NEXT: [[TMP1:%.*]] = fcmp one float [[TMP0]], 0x7FF0000000000000
+// NCRDIV-NEXT: [[CONV:%.*]] = zext i1 [[TMP1]] to i32
+// NCRDIV-NEXT: ret i32 [[CONV]]
+//
// AMDGCNSPIRV-LABEL: @test___finitef(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call addrspace(4) float @llvm.fabs.f32(float [[X:%.*]])
@@ -2058,6 +2400,13 @@ extern "C" __device__ BOOL_TYPE test___finitef(float x) {
// APPROX-NEXT: [[CONV:%.*]] = zext i1 [[TMP1]] to i32
// APPROX-NEXT: ret i32 [[CONV]]
//
+// NCRDIV-LABEL: @test___finite(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call double @llvm.fabs.f64(double [[X:%.*]])
+// NCRDIV-NEXT: [[TMP1:%.*]] = fcmp one double [[TMP0]], 0x7FF0000000000000
+// NCRDIV-NEXT: [[CONV:%.*]] = zext i1 [[TMP1]] to i32
+// NCRDIV-NEXT: ret i32 [[CONV]]
+//
// AMDGCNSPIRV-LABEL: @test___finite(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call addrspace(4) double @llvm.fabs.f64(double [[X:%.*]])
@@ -2087,6 +2436,13 @@ extern "C" __device__ BOOL_TYPE test___finite(double x) {
// APPROX-NEXT: [[CONV:%.*]] = zext i1 [[TMP1]] to i32
// APPROX-NEXT: ret i32 [[CONV]]
//
+// NCRDIV-LABEL: @test___isinff(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call float @llvm.fabs.f32(float [[X:%.*]])
+// NCRDIV-NEXT: [[TMP1:%.*]] = fcmp oeq float [[TMP0]], 0x7FF0000000000000
+// NCRDIV-NEXT: [[CONV:%.*]] = zext i1 [[TMP1]] to i32
+// NCRDIV-NEXT: ret i32 [[CONV]]
+//
// AMDGCNSPIRV-LABEL: @test___isinff(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call addrspace(4) float @llvm.fabs.f32(float [[X:%.*]])
@@ -2116,6 +2472,13 @@ extern "C" __device__ BOOL_TYPE test___isinff(float x) {
// APPROX-NEXT: [[CONV:%.*]] = zext i1 [[TMP1]] to i32
// APPROX-NEXT: ret i32 [[CONV]]
//
+// NCRDIV-LABEL: @test___isinf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call double @llvm.fabs.f64(double [[X:%.*]])
+// NCRDIV-NEXT: [[TMP1:%.*]] = fcmp oeq double [[TMP0]], 0x7FF0000000000000
+// NCRDIV-NEXT: [[CONV:%.*]] = zext i1 [[TMP1]] to i32
+// NCRDIV-NEXT: ret i32 [[CONV]]
+//
// AMDGCNSPIRV-LABEL: @test___isinf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call addrspace(4) double @llvm.fabs.f64(double [[X:%.*]])
@@ -2143,6 +2506,12 @@ extern "C" __device__ BOOL_TYPE test___isinf(double x) {
// APPROX-NEXT: [[CONV:%.*]] = zext i1 [[TMP0]] to i32
// APPROX-NEXT: ret i32 [[CONV]]
//
+// NCRDIV-LABEL: @test___isnanf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = fcmp uno float [[X:%.*]], 0.000000e+00
+// NCRDIV-NEXT: [[CONV:%.*]] = zext i1 [[TMP0]] to i32
+// NCRDIV-NEXT: ret i32 [[CONV]]
+//
// AMDGCNSPIRV-LABEL: @test___isnanf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = fcmp uno float [[X:%.*]], 0.000000e+00
@@ -2169,6 +2538,12 @@ extern "C" __device__ BOOL_TYPE test___isnanf(float x) {
// APPROX-NEXT: [[CONV:%.*]] = zext i1 [[TMP0]] to i32
// APPROX-NEXT: ret i32 [[CONV]]
//
+// NCRDIV-LABEL: @test___isnan(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = fcmp uno double [[X:%.*]], 0.000000e+00
+// NCRDIV-NEXT: [[CONV:%.*]] = zext i1 [[TMP0]] to i32
+// NCRDIV-NEXT: ret i32 [[CONV]]
+//
// AMDGCNSPIRV-LABEL: @test___isnan(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = fcmp uno double [[X:%.*]], 0.000000e+00
@@ -2194,6 +2569,11 @@ extern "C" __device__ BOOL_TYPE test___isnan(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_j0f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_j0f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -2218,6 +2598,11 @@ extern "C" __device__ float test_j0f(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_j0(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_j0(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -2242,6 +2627,11 @@ extern "C" __device__ double test_j0(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_j1f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_j1f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -2266,6 +2656,11 @@ extern "C" __device__ float test_j1f(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_j1(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_j1(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -2374,6 +2769,39 @@ extern "C" __device__ double test_j1(double x) {
// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// APPROX-NEXT: ret float [[RETVAL_0_I]]
//
+// NCRDIV-LABEL: @test_jnf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// NCRDIV-NEXT: i32 0, label [[IF_THEN_I:%.*]]
+// NCRDIV-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
+// NCRDIV-NEXT: ]
+// NCRDIV: if.then.i:
+// NCRDIV-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
+// NCRDIV: if.then2.i:
+// NCRDIV-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: br label [[_ZL3JNFIF_EXIT]]
+// NCRDIV: if.end4.i:
+// NCRDIV-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// NCRDIV-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// NCRDIV: for.body.i:
+// NCRDIV-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
+// NCRDIV-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
+// NCRDIV-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]], !fpmath [[META12]]
+// NCRDIV-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
+// NCRDIV-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
+// NCRDIV-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// NCRDIV-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// NCRDIV-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]]
+// NCRDIV: _ZL3jnfif.exit:
+// NCRDIV-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// NCRDIV-NEXT: ret float [[RETVAL_0_I]]
+//
// AMDGCNSPIRV-LABEL: @test_jnf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
@@ -2510,6 +2938,39 @@ extern "C" __device__ float test_jnf(int x, float y) {
// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// APPROX-NEXT: ret double [[RETVAL_0_I]]
//
+// NCRDIV-LABEL: @test_jn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// NCRDIV-NEXT: i32 0, label [[IF_THEN_I:%.*]]
+// NCRDIV-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
+// NCRDIV-NEXT: ]
+// NCRDIV: if.then.i:
+// NCRDIV-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
+// NCRDIV: if.then2.i:
+// NCRDIV-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: br label [[_ZL2JNID_EXIT]]
+// NCRDIV: if.end4.i:
+// NCRDIV-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// NCRDIV-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// NCRDIV: for.body.i:
+// NCRDIV-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
+// NCRDIV-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
+// NCRDIV-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
+// NCRDIV-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
+// NCRDIV-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
+// NCRDIV-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// NCRDIV-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// NCRDIV-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP16:![0-9]+]]
+// NCRDIV: _ZL2jnid.exit:
+// NCRDIV-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// NCRDIV-NEXT: ret double [[RETVAL_0_I]]
+//
// AMDGCNSPIRV-LABEL: @test_jn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
@@ -2562,6 +3023,11 @@ extern "C" __device__ double test_jn(int x, double y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_ldexpf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_ldexpf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
@@ -2586,6 +3052,11 @@ extern "C" __device__ float test_ldexpf(float x, int y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_ldexp(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_ldexp(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
@@ -2610,6 +3081,11 @@ extern "C" __device__ double test_ldexp(double x, int y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_lgammaf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_lgammaf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -2634,6 +3110,11 @@ extern "C" __device__ float test_lgammaf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_lgamma(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_lgamma(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -2661,6 +3142,12 @@ extern "C" __device__ double test_lgamma(double x) {
// APPROX-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// APPROX-NEXT: ret i64 [[CONV_I]]
//
+// NCRDIV-LABEL: @test_llrintf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
+// NCRDIV-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
+// NCRDIV-NEXT: ret i64 [[CONV_I]]
+//
// AMDGCNSPIRV-LABEL: @test_llrintf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract addrspace(4) float @llvm.rint.f32(float [[X:%.*]])
@@ -2689,6 +3176,12 @@ extern "C" __device__ long long int test_llrintf(float x) {
// APPROX-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// APPROX-NEXT: ret i64 [[CONV_I]]
//
+// NCRDIV-LABEL: @test_llrint(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
+// NCRDIV-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
+// NCRDIV-NEXT: ret i64 [[CONV_I]]
+//
// AMDGCNSPIRV-LABEL: @test_llrint(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract addrspace(4) double @llvm.rint.f64(double [[X:%.*]])
@@ -2717,6 +3210,12 @@ extern "C" __device__ long long int test_llrint(double x) {
// APPROX-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// APPROX-NEXT: ret i64 [[CONV_I]]
//
+// NCRDIV-LABEL: @test_llroundf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]])
+// NCRDIV-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
+// NCRDIV-NEXT: ret i64 [[CONV_I]]
+//
// AMDGCNSPIRV-LABEL: @test_llroundf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract addrspace(4) float @llvm.round.f32(float [[X:%.*]])
@@ -2745,6 +3244,12 @@ extern "C" __device__ long long int test_llroundf(float x) {
// APPROX-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// APPROX-NEXT: ret i64 [[CONV_I]]
//
+// NCRDIV-LABEL: @test_llround(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]])
+// NCRDIV-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
+// NCRDIV-NEXT: ret i64 [[CONV_I]]
+//
// AMDGCNSPIRV-LABEL: @test_llround(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract addrspace(4) double @llvm.round.f64(double [[X:%.*]])
@@ -2770,6 +3275,11 @@ extern "C" __device__ long long int test_llround(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_log10f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_log10f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log10.f32(float [[X:%.*]])
@@ -2794,6 +3304,11 @@ extern "C" __device__ float test_log10f(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_log10(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_log10(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -2818,6 +3333,11 @@ extern "C" __device__ double test_log10(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_log1pf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_log1pf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -2842,6 +3362,11 @@ extern "C" __device__ float test_log1pf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_log1p(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_log1p(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -2866,6 +3391,11 @@ extern "C" __device__ double test_log1p(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.amdgcn.log.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_log2f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log2.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_log2f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log2.f32(float [[X:%.*]])
@@ -2890,6 +3420,11 @@ extern "C" __device__ float test_log2f(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_log2(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_log2(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -2914,6 +3449,11 @@ extern "C" __device__ double test_log2(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_logbf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_logbf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR12]]
@@ -2938,6 +3478,11 @@ extern "C" __device__ float test_logbf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_logb(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_logb(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR12]]
@@ -2962,6 +3507,11 @@ extern "C" __device__ double test_logb(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_logf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_logf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log.f32(float [[X:%.*]])
@@ -2989,6 +3539,12 @@ extern "C" __device__ float test_logf(float x) {
// APPROX-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// APPROX-NEXT: ret i64 [[CONV_I]]
//
+// NCRDIV-LABEL: @test_lrintf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
+// NCRDIV-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
+// NCRDIV-NEXT: ret i64 [[CONV_I]]
+//
// AMDGCNSPIRV-LABEL: @test_lrintf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract addrspace(4) float @llvm.rint.f32(float [[X:%.*]])
@@ -3017,6 +3573,12 @@ extern "C" __device__ long int test_lrintf(float x) {
// APPROX-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// APPROX-NEXT: ret i64 [[CONV_I]]
//
+// NCRDIV-LABEL: @test_lrint(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
+// NCRDIV-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
+// NCRDIV-NEXT: ret i64 [[CONV_I]]
+//
// AMDGCNSPIRV-LABEL: @test_lrint(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract addrspace(4) double @llvm.rint.f64(double [[X:%.*]])
@@ -3045,6 +3607,12 @@ extern "C" __device__ long int test_lrint(double x) {
// APPROX-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// APPROX-NEXT: ret i64 [[CONV_I]]
//
+// NCRDIV-LABEL: @test_lroundf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]])
+// NCRDIV-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
+// NCRDIV-NEXT: ret i64 [[CONV_I]]
+//
// AMDGCNSPIRV-LABEL: @test_lroundf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract addrspace(4) float @llvm.round.f32(float [[X:%.*]])
@@ -3073,6 +3641,12 @@ extern "C" __device__ long int test_lroundf(float x) {
// APPROX-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// APPROX-NEXT: ret i64 [[CONV_I]]
//
+// NCRDIV-LABEL: @test_lround(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]])
+// NCRDIV-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
+// NCRDIV-NEXT: ret i64 [[CONV_I]]
+//
// AMDGCNSPIRV-LABEL: @test_lround(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract addrspace(4) double @llvm.round.f64(double [[X:%.*]])
@@ -3113,6 +3687,16 @@ extern "C" __device__ long int test_lround(double x) {
// APPROX-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_modff(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
+// NCRDIV-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15:[0-9]+]]
+// NCRDIV-NEXT: [[CALL_I:%.*]] = call contract noundef float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA17:![0-9]+]]
+// NCRDIV-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA17]]
+// NCRDIV-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_modff(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4
@@ -3158,6 +3742,16 @@ extern "C" __device__ float test_modff(float x, float* y) {
// APPROX-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_modf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// NCRDIV-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: [[CALL_I:%.*]] = call contract noundef double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA19:![0-9]+]]
+// NCRDIV-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA19]]
+// NCRDIV-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_modf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8
@@ -3367,6 +3961,101 @@ extern "C" __device__ double test_modf(double x, double* y) {
// APPROX-NEXT: [[TMP10:%.*]] = bitcast i32 [[BF_SET9_I]] to float
// APPROX-NEXT: ret float [[TMP10]]
//
+// NCRDIV-LABEL: @test_nanf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA4]]
+// NCRDIV-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
+// NCRDIV-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I14_I_I:%.*]]
+// NCRDIV: if.then.i.i:
+// NCRDIV-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
+// NCRDIV-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA4]]
+// NCRDIV-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
+// NCRDIV-NEXT: i8 120, label [[WHILE_COND_I30_I_I_PREHEADER:%.*]]
+// NCRDIV-NEXT: i8 88, label [[WHILE_COND_I30_I_I_PREHEADER]]
+// NCRDIV-NEXT: ]
+// NCRDIV: while.cond.i30.i.i.preheader:
+// NCRDIV-NEXT: br label [[WHILE_COND_I30_I_I:%.*]]
+// NCRDIV: while.cond.i30.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I30_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I36_I_I]] ], [ 0, [[WHILE_COND_I30_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[TBAA4]]
+// NCRDIV-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]], label [[WHILE_BODY_I34_I_I:%.*]]
+// NCRDIV: while.body.i34.i.i:
+// NCRDIV-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
+// NCRDIV-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// NCRDIV-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
+// NCRDIV: if.else.i.i.i:
+// NCRDIV-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
+// NCRDIV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// NCRDIV-NEXT: br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label [[IF_ELSE17_I_I_I:%.*]]
+// NCRDIV: if.else17.i.i.i:
+// NCRDIV-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
+// NCRDIV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
+// NCRDIV-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I36_I_I]]
+// NCRDIV: if.end31.i.i.i:
+// NCRDIV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
+// NCRDIV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// NCRDIV-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
+// NCRDIV-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1
+// NCRDIV-NEXT: br label [[CLEANUP_I36_I_I]]
+// NCRDIV: cleanup.i36.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], [[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], [[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I30_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP11]]
+// NCRDIV: while.cond.i.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
+// NCRDIV-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// NCRDIV-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA4]]
+// NCRDIV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I_I_I:%.*]]
+// NCRDIV: while.body.i.i.i:
+// NCRDIV-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// NCRDIV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// NCRDIV-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label [[CLEANUP_I_I_I]]
+// NCRDIV: if.then.i.i.i:
+// NCRDIV-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
+// NCRDIV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// NCRDIV-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
+// NCRDIV-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
+// NCRDIV-NEXT: br label [[CLEANUP_I_I_I]]
+// NCRDIV: cleanup.i.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// NCRDIV-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// NCRDIV-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP7]]
+// NCRDIV: while.cond.i14.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
+// NCRDIV-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], [[CLEANUP_I20_I_I]] ], [ 0, [[ENTRY]] ]
+// NCRDIV-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[TBAA4]]
+// NCRDIV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I17_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I18_I_I:%.*]]
+// NCRDIV: while.body.i18.i.i:
+// NCRDIV-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// NCRDIV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// NCRDIV-NEXT: br i1 [[OR_COND_I19_I_I]], label [[IF_THEN_I24_I_I:%.*]], label [[CLEANUP_I20_I_I]]
+// NCRDIV: if.then.i24.i.i:
+// NCRDIV-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
+// NCRDIV-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
+// NCRDIV-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
+// NCRDIV-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]], [[CONV5_I26_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1
+// NCRDIV-NEXT: br label [[CLEANUP_I20_I_I]]
+// NCRDIV: cleanup.i20.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [ [[INCDEC_PTR_I29_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]], [[WHILE_BODY_I18_I_I]] ]
+// NCRDIV-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_BODY_I18_I_I]] ]
+// NCRDIV-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]]
+// NCRDIV: _ZL4nanfPKc.exit:
+// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [ [[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ], [ 0, [[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ]
+// NCRDIV-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
+// NCRDIV-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
+// NCRDIV-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]], 2143289344
+// NCRDIV-NEXT: [[TMP10:%.*]] = bitcast i32 [[BF_SET9_I]] to float
+// NCRDIV-NEXT: ret float [[TMP10]]
+//
// AMDGCNSPIRV-LABEL: @test_nanf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load i8, ptr addrspace(4) [[TAG:%.*]], align 1, !tbaa [[TBAA5]]
@@ -3650,6 +4339,100 @@ extern "C" __device__ float test_nanf(const char *tag) {
// APPROX-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
// APPROX-NEXT: ret double [[TMP10]]
//
+// NCRDIV-LABEL: @test_nan(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA4]]
+// NCRDIV-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
+// NCRDIV-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I14_I_I:%.*]]
+// NCRDIV: if.then.i.i:
+// NCRDIV-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
+// NCRDIV-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA4]]
+// NCRDIV-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
+// NCRDIV-NEXT: i8 120, label [[WHILE_COND_I30_I_I_PREHEADER:%.*]]
+// NCRDIV-NEXT: i8 88, label [[WHILE_COND_I30_I_I_PREHEADER]]
+// NCRDIV-NEXT: ]
+// NCRDIV: while.cond.i30.i.i.preheader:
+// NCRDIV-NEXT: br label [[WHILE_COND_I30_I_I:%.*]]
+// NCRDIV: while.cond.i30.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I30_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I36_I_I]] ], [ 0, [[WHILE_COND_I30_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[TBAA4]]
+// NCRDIV-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL3NANPKC_EXIT:%.*]], label [[WHILE_BODY_I34_I_I:%.*]]
+// NCRDIV: while.body.i34.i.i:
+// NCRDIV-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
+// NCRDIV-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// NCRDIV-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
+// NCRDIV: if.else.i.i.i:
+// NCRDIV-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
+// NCRDIV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// NCRDIV-NEXT: br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label [[IF_ELSE17_I_I_I:%.*]]
+// NCRDIV: if.else17.i.i.i:
+// NCRDIV-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
+// NCRDIV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
+// NCRDIV-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I36_I_I]]
+// NCRDIV: if.end31.i.i.i:
+// NCRDIV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
+// NCRDIV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// NCRDIV-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
+// NCRDIV-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1
+// NCRDIV-NEXT: br label [[CLEANUP_I36_I_I]]
+// NCRDIV: cleanup.i36.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], [[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], [[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I30_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP11]]
+// NCRDIV: while.cond.i.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
+// NCRDIV-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// NCRDIV-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA4]]
+// NCRDIV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I_I_I:%.*]]
+// NCRDIV: while.body.i.i.i:
+// NCRDIV-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// NCRDIV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// NCRDIV-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label [[CLEANUP_I_I_I]]
+// NCRDIV: if.then.i.i.i:
+// NCRDIV-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
+// NCRDIV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// NCRDIV-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
+// NCRDIV-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
+// NCRDIV-NEXT: br label [[CLEANUP_I_I_I]]
+// NCRDIV: cleanup.i.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// NCRDIV-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// NCRDIV-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP7]]
+// NCRDIV: while.cond.i14.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
+// NCRDIV-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], [[CLEANUP_I20_I_I]] ], [ 0, [[ENTRY]] ]
+// NCRDIV-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[TBAA4]]
+// NCRDIV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I17_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I18_I_I:%.*]]
+// NCRDIV: while.body.i18.i.i:
+// NCRDIV-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// NCRDIV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// NCRDIV-NEXT: br i1 [[OR_COND_I19_I_I]], label [[IF_THEN_I24_I_I:%.*]], label [[CLEANUP_I20_I_I]]
+// NCRDIV: if.then.i24.i.i:
+// NCRDIV-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
+// NCRDIV-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
+// NCRDIV-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
+// NCRDIV-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]], [[CONV5_I26_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1
+// NCRDIV-NEXT: br label [[CLEANUP_I20_I_I]]
+// NCRDIV: cleanup.i20.i.i:
+// NCRDIV-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [ [[INCDEC_PTR_I29_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]], [[WHILE_BODY_I18_I_I]] ]
+// NCRDIV-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_BODY_I18_I_I]] ]
+// NCRDIV-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]]
+// NCRDIV: _ZL3nanPKc.exit:
+// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [ [[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ], [ 0, [[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ]
+// NCRDIV-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247
+// NCRDIV-NEXT: [[BF_SET9_I:%.*]] = or disjoint i64 [[BF_VALUE_I]], 9221120237041090560
+// NCRDIV-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
+// NCRDIV-NEXT: ret double [[TMP10]]
+//
// AMDGCNSPIRV-LABEL: @test_nan(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load i8, ptr addrspace(4) [[TAG:%.*]], align 1, !tbaa [[TBAA5]]
@@ -3752,6 +4535,10 @@ extern "C" __device__ double test_nan(const char *tag) {
// APPROX-NEXT: entry:
// APPROX-NEXT: ret float 0x7FF8000000000000
//
+// NCRDIV-LABEL: @test_nanf_emptystr(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: ret float 0x7FF8000000000000
+//
// AMDGCNSPIRV-LABEL: @test_nanf_emptystr(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: ret float 0x7FF8000000000000
@@ -3772,6 +4559,10 @@ extern "C" __device__ float test_nanf_emptystr() {
// APPROX-NEXT: entry:
// APPROX-NEXT: ret double 0x7FF8000000000000
//
+// NCRDIV-LABEL: @test_nan_emptystr(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: ret double 0x7FF8000000000000
+//
// AMDGCNSPIRV-LABEL: @test_nan_emptystr(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: ret double 0x7FF8000000000000
@@ -3792,6 +4583,10 @@ extern "C" __device__ double test_nan_emptystr() {
// APPROX-NEXT: entry:
// APPROX-NEXT: ret float 0x7FF8000000000000
//
+// NCRDIV-LABEL: @test_nanf_fill(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: ret float 0x7FF8000000000000
+//
// AMDGCNSPIRV-LABEL: @test_nanf_fill(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: ret float 0x7FF8000000000000
@@ -3812,6 +4607,10 @@ extern "C" __device__ float test_nanf_fill() {
// APPROX-NEXT: entry:
// APPROX-NEXT: ret double 0x7FF8000000000000
//
+// NCRDIV-LABEL: @test_nan_fill(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: ret double 0x7FF8000000000000
+//
// AMDGCNSPIRV-LABEL: @test_nan_fill(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: ret double 0x7FF8000000000000
@@ -3835,6 +4634,11 @@ extern "C" __device__ double test_nan_fill() {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.nearbyint.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_nearbyintf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.nearbyint.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_nearbyintf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.nearbyint.f32(float [[X:%.*]])
@@ -3859,6 +4663,11 @@ extern "C" __device__ float test_nearbyintf(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.nearbyint.f64(double [[X:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_nearbyint(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.nearbyint.f64(double [[X:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_nearbyint(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.nearbyint.f64(double [[X:%.*]])
@@ -3883,6 +4692,11 @@ extern "C" __device__ double test_nearbyint(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_nextafterf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_nextafterf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
@@ -3907,6 +4721,11 @@ extern "C" __device__ float test_nextafterf(float x, float y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_nextafter(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_nextafter(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
@@ -3931,6 +4750,11 @@ extern "C" __device__ double test_nextafter(double x, double y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_norm3df(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_norm3df(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR12]]
@@ -3955,6 +4779,11 @@ extern "C" __device__ float test_norm3df(float x, float y, float z) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_norm3d(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_norm3d(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR12]]
@@ -3979,6 +4808,11 @@ extern "C" __device__ double test_norm3d(double x, double y, double z) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_norm4df(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_norm4df(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR12]]
@@ -4003,6 +4837,11 @@ extern "C" __device__ float test_norm4df(float x, float y, float z, float w) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_norm4d(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_norm4d(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR12]]
@@ -4027,6 +4866,11 @@ extern "C" __device__ double test_norm4d(double x, double y, double z, double w)
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_normcdff(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_normcdff(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -4051,6 +4895,11 @@ extern "C" __device__ float test_normcdff(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_normcdf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_normcdf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -4075,6 +4924,11 @@ extern "C" __device__ double test_normcdf(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_normcdfinvf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_normcdfinvf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -4099,6 +4953,11 @@ extern "C" __device__ float test_normcdfinvf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_normcdfinv(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_normcdfinv(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -4168,6 +5027,26 @@ extern "C" __device__ double test_normcdfinv(double x) {
// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]])
// APPROX-NEXT: ret float [[TMP1]]
//
+// NCRDIV-LABEL: @test_normf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// NCRDIV: while.body.i:
+// NCRDIV-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// NCRDIV-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// NCRDIV-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// NCRDIV-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// NCRDIV-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA17]]
+// NCRDIV-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
+// NCRDIV-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 4
+// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
+// NCRDIV: _ZL5normfiPKf.exit:
+// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]]), !fpmath [[META22:![0-9]+]]
+// NCRDIV-NEXT: ret float [[TMP1]]
+//
// AMDGCNSPIRV-LABEL: @test_normf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
@@ -4252,6 +5131,26 @@ extern "C" __device__ float test_normf(int x, const float *y) {
// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]])
// APPROX-NEXT: ret double [[TMP1]]
//
+// NCRDIV-LABEL: @test_norm(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// NCRDIV: while.body.i:
+// NCRDIV-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// NCRDIV-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// NCRDIV-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// NCRDIV-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// NCRDIV-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA19]]
+// NCRDIV-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
+// NCRDIV-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 8
+// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
+// NCRDIV: _ZL4normiPKd.exit:
+// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]])
+// NCRDIV-NEXT: ret double [[TMP1]]
+//
// AMDGCNSPIRV-LABEL: @test_norm(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
@@ -4291,6 +5190,11 @@ extern "C" __device__ double test_norm(int x, const double *y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_powf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_powf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
@@ -4315,6 +5219,11 @@ extern "C" __device__ float test_powf(float x, float y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_pow(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_pow(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
@@ -4339,6 +5248,11 @@ extern "C" __device__ double test_pow(double x, double y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_powif(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_powif(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
@@ -4363,6 +5277,11 @@ extern "C" __device__ float test_powif(float x, int y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_powi(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_powi(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
@@ -4387,6 +5306,11 @@ extern "C" __device__ double test_powi(double x, int y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rcbrtf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rcbrtf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -4411,6 +5335,11 @@ extern "C" __device__ float test_rcbrtf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rcbrt(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rcbrt(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -4435,6 +5364,11 @@ extern "C" __device__ double test_rcbrt(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_remainderf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_remainderf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
@@ -4459,6 +5393,11 @@ extern "C" __device__ float test_remainderf(float x, float y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_remainder(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_remainder(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
@@ -4498,6 +5437,16 @@ extern "C" __device__ double test_remainder(double x, double y) {
// APPROX-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_remquof(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
+// NCRDIV-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: [[CALL_I:%.*]] = call contract noundef float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]]
+// NCRDIV-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
+// NCRDIV-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_remquof(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4
@@ -4543,6 +5492,16 @@ extern "C" __device__ float test_remquof(float x, float y, int* z) {
// APPROX-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_remquo(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
+// NCRDIV-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: [[CALL_I:%.*]] = call contract noundef double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]]
+// NCRDIV-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
+// NCRDIV-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_remquo(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4
@@ -4573,6 +5532,11 @@ extern "C" __device__ double test_remquo(double x, double y, int* z) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rhypotf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rhypotf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR12]]
@@ -4597,6 +5561,11 @@ extern "C" __device__ float test_rhypotf(float x, float y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rhypot(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rhypot(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR12]]
@@ -4621,6 +5590,11 @@ extern "C" __device__ double test_rhypot(double x, double y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.rint.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_rintf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.rint.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_rintf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.rint.f32(float [[X:%.*]])
@@ -4645,6 +5619,11 @@ extern "C" __device__ float test_rintf(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.rint.f64(double [[X:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_rint(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.rint.f64(double [[X:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_rint(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.rint.f64(double [[X:%.*]])
@@ -4714,6 +5693,26 @@ extern "C" __device__ double test_rint(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rnormf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// NCRDIV: while.body.i:
+// NCRDIV-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// NCRDIV-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// NCRDIV-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// NCRDIV-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// NCRDIV-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA17]]
+// NCRDIV-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
+// NCRDIV-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 4
+// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// NCRDIV: _ZL6rnormfiPKf.exit:
+// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rnormf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
@@ -4798,6 +5797,26 @@ extern "C" __device__ float test_rnormf(int x, const float* y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rnorm(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// NCRDIV: while.body.i:
+// NCRDIV-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// NCRDIV-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// NCRDIV-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// NCRDIV-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// NCRDIV-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA19]]
+// NCRDIV-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
+// NCRDIV-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 8
+// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]]
+// NCRDIV: _ZL5rnormiPKd.exit:
+// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rnorm(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
@@ -4837,6 +5856,11 @@ extern "C" __device__ double test_rnorm(int x, const double* y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rnorm3df(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rnorm3df(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR12]]
@@ -4861,6 +5885,11 @@ extern "C" __device__ float test_rnorm3df(float x, float y, float z) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rnorm3d(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rnorm3d(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR12]]
@@ -4885,6 +5914,11 @@ extern "C" __device__ double test_rnorm3d(double x, double y, double z) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rnorm4df(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rnorm4df(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR12]]
@@ -4909,6 +5943,11 @@ extern "C" __device__ float test_rnorm4df(float x, float y, float z, float w) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rnorm4d(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rnorm4d(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR12]]
@@ -4933,6 +5972,11 @@ extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.round.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_roundf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.round.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_roundf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.round.f32(float [[X:%.*]])
@@ -4957,6 +6001,11 @@ extern "C" __device__ float test_roundf(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.round.f64(double [[X:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_round(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.round.f64(double [[X:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_round(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.round.f64(double [[X:%.*]])
@@ -4981,6 +6030,11 @@ extern "C" __device__ double test_round(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rsqrtf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rsqrtf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -5005,6 +6059,11 @@ extern "C" __device__ float test_rsqrtf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_rsqrt(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_rsqrt(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -5035,6 +6094,13 @@ extern "C" __device__ double test_rsqrt(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[CONV_I]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_scalblnf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[SPEC_STORE_SELECT_I:%.*]] = tail call i64 @llvm.smax.i64(i64 [[Y:%.*]], i64 -2147483648)
+// NCRDIV-NEXT: [[CONV_I:%.*]] = trunc i64 [[SPEC_STORE_SELECT_I]] to i32
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[CONV_I]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_scalblnf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[SPEC_STORE_SELECT_I:%.*]] = tail call addrspace(4) i64 @llvm.smax.i64(i64 [[Y:%.*]], i64 -2147483648)
@@ -5067,6 +6133,13 @@ extern "C" __device__ float test_scalblnf(float x, long int y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[CONV_I]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_scalbln(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[SPEC_STORE_SELECT_I:%.*]] = tail call i64 @llvm.smax.i64(i64 [[Y:%.*]], i64 -2147483648)
+// NCRDIV-NEXT: [[CONV_I:%.*]] = trunc i64 [[SPEC_STORE_SELECT_I]] to i32
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[CONV_I]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_scalbln(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[SPEC_STORE_SELECT_I:%.*]] = tail call addrspace(4) i64 @llvm.smax.i64(i64 [[Y:%.*]], i64 -2147483648)
@@ -5093,6 +6166,11 @@ extern "C" __device__ double test_scalbln(double x, long int y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_scalbnf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_scalbnf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
@@ -5117,6 +6195,11 @@ extern "C" __device__ float test_scalbnf(float x, int y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_scalbn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_scalbn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
@@ -5193,6 +6276,17 @@ extern "C" __device__ BOOL_TYPE test___signbit(double x) {
// APPROX-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
// APPROX-NEXT: ret void
//
+// NCRDIV-LABEL: @test_sincosf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
+// NCRDIV-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR14]]
+// NCRDIV-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA17]]
+// NCRDIV-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA17]]
+// NCRDIV-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA17]]
+// NCRDIV-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: ret void
+//
// AMDGCNSPIRV-LABEL: @test_sincosf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4
@@ -5242,6 +6336,17 @@ extern "C" __device__ void test_sincosf(float x, float *y, float *z) {
// APPROX-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
// APPROX-NEXT: ret void
//
+// NCRDIV-LABEL: @test_sincos(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// NCRDIV-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR14]]
+// NCRDIV-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA19]]
+// NCRDIV-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA19]]
+// NCRDIV-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA19]]
+// NCRDIV-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: ret void
+//
// AMDGCNSPIRV-LABEL: @test_sincos(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8
@@ -5291,6 +6396,17 @@ extern "C" __device__ void test_sincos(double x, double *y, double *z) {
// APPROX-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
// APPROX-NEXT: ret void
//
+// NCRDIV-LABEL: @test_sincospif(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
+// NCRDIV-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR14]]
+// NCRDIV-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA17]]
+// NCRDIV-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA17]]
+// NCRDIV-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA17]]
+// NCRDIV-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: ret void
+//
// AMDGCNSPIRV-LABEL: @test_sincospif(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca float, align 4
@@ -5340,6 +6456,17 @@ extern "C" __device__ void test_sincospif(float x, float *y, float *z) {
// APPROX-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
// APPROX-NEXT: ret void
//
+// NCRDIV-LABEL: @test_sincospi(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// NCRDIV-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR14]]
+// NCRDIV-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA19]]
+// NCRDIV-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA19]]
+// NCRDIV-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA19]]
+// NCRDIV-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR15]]
+// NCRDIV-NEXT: ret void
+//
// AMDGCNSPIRV-LABEL: @test_sincospi(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[__TMP_I:%.*]] = alloca double, align 8
@@ -5371,6 +6498,11 @@ extern "C" __device__ void test_sincospi(double x, double *y, double *z) {
// APPROX-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I1]]
//
+// NCRDIV-LABEL: @test_sinf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_sinf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -5395,6 +6527,11 @@ extern "C" __device__ float test_sinf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_sin(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_sin(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -5419,6 +6556,11 @@ extern "C" __device__ double test_sin(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_sinpif(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_sinpif(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -5443,6 +6585,11 @@ extern "C" __device__ float test_sinpif(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_sinpi(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_sinpi(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -5467,6 +6614,11 @@ extern "C" __device__ double test_sinpi(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_sqrtf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X:%.*]]), !fpmath [[META22]]
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_sqrtf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.sqrt.f32(float [[X:%.*]])
@@ -5491,6 +6643,11 @@ extern "C" __device__ float test_sqrtf(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[X:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_sqrt(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[X:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_sqrt(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.sqrt.f64(double [[X:%.*]])
@@ -5515,6 +6672,11 @@ extern "C" __device__ double test_sqrt(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_tanf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_tanf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -5539,6 +6701,11 @@ extern "C" __device__ float test_tanf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_tan(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_tan(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -5563,6 +6730,11 @@ extern "C" __device__ double test_tan(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_tanhf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_tanhf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR13]]
@@ -5587,6 +6759,11 @@ extern "C" __device__ float test_tanhf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_tanh(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_tanh(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR13]]
@@ -5611,6 +6788,11 @@ extern "C" __device__ double test_tanh(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_tgammaf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_tgammaf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -5635,6 +6817,11 @@ extern "C" __device__ float test_tgammaf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_tgamma(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_tgamma(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -5659,6 +6846,11 @@ extern "C" __device__ double test_tgamma(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.trunc.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_truncf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.trunc.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_truncf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.trunc.f32(float [[X:%.*]])
@@ -5683,6 +6875,11 @@ extern "C" __device__ float test_truncf(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.trunc.f64(double [[X:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_trunc(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.trunc.f64(double [[X:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_trunc(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.trunc.f64(double [[X:%.*]])
@@ -5707,6 +6904,11 @@ extern "C" __device__ double test_trunc(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_y0f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_y0f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -5731,6 +6933,11 @@ extern "C" __device__ float test_y0f(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_y0(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_y0(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -5755,6 +6962,11 @@ extern "C" __device__ double test_y0(double x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test_y1f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_y1f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -5779,6 +6991,11 @@ extern "C" __device__ float test_y1f(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret double [[CALL_I]]
//
+// NCRDIV-LABEL: @test_y1(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret double [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test_y1(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -5887,6 +7104,39 @@ extern "C" __device__ double test_y1(double x) {
// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// APPROX-NEXT: ret float [[RETVAL_0_I]]
//
+// NCRDIV-LABEL: @test_ynf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// NCRDIV-NEXT: i32 0, label [[IF_THEN_I:%.*]]
+// NCRDIV-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
+// NCRDIV-NEXT: ]
+// NCRDIV: if.then.i:
+// NCRDIV-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
+// NCRDIV: if.then2.i:
+// NCRDIV-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: br label [[_ZL3YNFIF_EXIT]]
+// NCRDIV: if.end4.i:
+// NCRDIV-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// NCRDIV-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// NCRDIV: for.body.i:
+// NCRDIV-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
+// NCRDIV-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
+// NCRDIV-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]], !fpmath [[META12]]
+// NCRDIV-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
+// NCRDIV-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
+// NCRDIV-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// NCRDIV-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// NCRDIV-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP26:![0-9]+]]
+// NCRDIV: _ZL3ynfif.exit:
+// NCRDIV-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// NCRDIV-NEXT: ret float [[RETVAL_0_I]]
+//
// AMDGCNSPIRV-LABEL: @test_ynf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
@@ -6023,6 +7273,39 @@ extern "C" __device__ float test_ynf(int x, float y) {
// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// APPROX-NEXT: ret double [[RETVAL_0_I]]
//
+// NCRDIV-LABEL: @test_yn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// NCRDIV-NEXT: i32 0, label [[IF_THEN_I:%.*]]
+// NCRDIV-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
+// NCRDIV-NEXT: ]
+// NCRDIV: if.then.i:
+// NCRDIV-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
+// NCRDIV: if.then2.i:
+// NCRDIV-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: br label [[_ZL2YNID_EXIT]]
+// NCRDIV: if.end4.i:
+// NCRDIV-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// NCRDIV-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// NCRDIV: for.body.i:
+// NCRDIV-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// NCRDIV-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
+// NCRDIV-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
+// NCRDIV-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
+// NCRDIV-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
+// NCRDIV-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
+// NCRDIV-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// NCRDIV-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// NCRDIV-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP27:![0-9]+]]
+// NCRDIV: _ZL2ynid.exit:
+// NCRDIV-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// NCRDIV-NEXT: ret double [[RETVAL_0_I]]
+//
// AMDGCNSPIRV-LABEL: @test_yn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
@@ -6075,6 +7358,11 @@ extern "C" __device__ double test_yn(int x, double y) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test___cosf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// 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_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -6102,6 +7390,12 @@ extern "C" __device__ float test___cosf(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test___exp10f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[MUL_I:%.*]] = fmul contract float [[X:%.*]], 0x400A934F00000000
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test___exp10f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[MUL_I:%.*]] = fmul contract float [[X:%.*]], 0x400A934F00000000
@@ -6130,6 +7424,12 @@ extern "C" __device__ float test___exp10f(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test___expf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[MUL_I:%.*]] = fmul contract float [[X:%.*]], 0x3FF7154760000000
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test___expf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[MUL_I:%.*]] = fmul contract float [[X:%.*]], 0x3FF7154760000000
@@ -6155,6 +7455,11 @@ extern "C" __device__ float test___expf(float x) {
// APPROX-NEXT: [[ADD_I:%.*]] = fadd contract float [[X:%.*]], [[Y:%.*]]
// APPROX-NEXT: ret float [[ADD_I]]
//
+// NCRDIV-LABEL: @test___fadd_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[ADD_I:%.*]] = fadd contract float [[X:%.*]], [[Y:%.*]]
+// NCRDIV-NEXT: ret float [[ADD_I]]
+//
// AMDGCNSPIRV-LABEL: @test___fadd_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[ADD_I:%.*]] = fadd contract float [[X:%.*]], [[Y:%.*]]
@@ -6179,6 +7484,11 @@ extern "C" __device__ float test___fadd_rn(float x, float y) {
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]]
// APPROX-NEXT: ret float [[DIV_I]]
//
+// NCRDIV-LABEL: @test___fdividef(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]], !fpmath [[META12]]
+// NCRDIV-NEXT: ret float [[DIV_I]]
+//
// AMDGCNSPIRV-LABEL: @test___fdividef(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]]
@@ -6203,6 +7513,11 @@ extern "C" __device__ float test___fdividef(float x, float y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test__fmaf_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test__fmaf_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
@@ -6227,6 +7542,11 @@ extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
// APPROX-NEXT: [[MUL_I:%.*]] = fmul contract float [[X:%.*]], [[Y:%.*]]
// APPROX-NEXT: ret float [[MUL_I]]
//
+// NCRDIV-LABEL: @test___fmul_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[MUL_I:%.*]] = fmul contract float [[X:%.*]], [[Y:%.*]]
+// NCRDIV-NEXT: ret float [[MUL_I]]
+//
// AMDGCNSPIRV-LABEL: @test___fmul_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[MUL_I:%.*]] = fmul contract float [[X:%.*]], [[Y:%.*]]
@@ -6251,6 +7571,11 @@ extern "C" __device__ float test___fmul_rn(float x, float y) {
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract float 1.000000e+00, [[X:%.*]]
// APPROX-NEXT: ret float [[DIV_I]]
//
+// NCRDIV-LABEL: @test___frcp_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[DIV_I:%.*]] = fdiv contract float 1.000000e+00, [[X:%.*]], !fpmath [[META12]]
+// NCRDIV-NEXT: ret float [[DIV_I]]
+//
// AMDGCNSPIRV-LABEL: @test___frcp_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[DIV_I:%.*]] = fdiv contract float 1.000000e+00, [[X:%.*]]
@@ -6275,6 +7600,11 @@ extern "C" __device__ float test___frcp_rn(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.amdgcn.rsq.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test___frsqrt_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.amdgcn.rsq.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test___frsqrt_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.amdgcn.rsq.f32(float [[X:%.*]])
@@ -6299,6 +7629,11 @@ extern "C" __device__ float test___frsqrt_rn(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR12]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test___fsqrt_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR12]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test___fsqrt_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR12]]
@@ -6323,6 +7658,11 @@ extern "C" __device__ float test___fsqrt_rn(float x) {
// APPROX-NEXT: [[SUB_I:%.*]] = fsub contract float [[X:%.*]], [[Y:%.*]]
// APPROX-NEXT: ret float [[SUB_I]]
//
+// NCRDIV-LABEL: @test___fsub_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[SUB_I:%.*]] = fsub contract float [[X:%.*]], [[Y:%.*]]
+// NCRDIV-NEXT: ret float [[SUB_I]]
+//
// AMDGCNSPIRV-LABEL: @test___fsub_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[SUB_I:%.*]] = fsub contract float [[X:%.*]], [[Y:%.*]]
@@ -6347,6 +7687,11 @@ extern "C" __device__ float test___fsub_rn(float x, float y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test___log10f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test___log10f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log10.f32(float [[X:%.*]])
@@ -6371,6 +7716,11 @@ extern "C" __device__ float test___log10f(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.amdgcn.log.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test___log2f(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.amdgcn.log.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test___log2f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.amdgcn.log.f32(float [[X:%.*]])
@@ -6395,6 +7745,11 @@ extern "C" __device__ float test___log2f(float x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test___logf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test___logf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log.f32(float [[X:%.*]])
@@ -6419,6 +7774,11 @@ extern "C" __device__ float test___logf(float x) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test___powf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test___powf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
@@ -6452,6 +7812,14 @@ extern "C" __device__ float test___powf(float x, float y) {
// APPROX-NEXT: [[COND5_I:%.*]] = select contract i1 [[CMP_I]], float 0.000000e+00, float [[COND_I]]
// APPROX-NEXT: ret float [[COND5_I]]
//
+// NCRDIV-LABEL: @test___saturatef(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CMP_I:%.*]] = fcmp contract olt float [[X:%.*]], 0.000000e+00
+// NCRDIV-NEXT: [[CMP1_I:%.*]] = fcmp contract ogt float [[X]], 1.000000e+00
+// NCRDIV-NEXT: [[COND_I:%.*]] = select contract i1 [[CMP1_I]], float 1.000000e+00, float [[X]]
+// NCRDIV-NEXT: [[COND5_I:%.*]] = select contract i1 [[CMP_I]], float 0.000000e+00, float [[COND_I]]
+// NCRDIV-NEXT: ret float [[COND5_I]]
+//
// AMDGCNSPIRV-LABEL: @test___saturatef(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CMP_I:%.*]] = fcmp contract olt float [[X:%.*]], 0.000000e+00
@@ -6488,6 +7856,14 @@ extern "C" __device__ float test___saturatef(float x) {
// APPROX-NEXT: store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA16]]
// APPROX-NEXT: ret void
//
+// NCRDIV-LABEL: @test___sincosf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA17]]
+// NCRDIV-NEXT: [[CALL1_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR14]]
+// NCRDIV-NEXT: store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA17]]
+// NCRDIV-NEXT: ret void
+//
// AMDGCNSPIRV-LABEL: @test___sincosf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func addrspace(4) float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -6515,6 +7891,11 @@ extern "C" __device__ void test___sincosf(float x, float *y, float *z) {
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// APPROX-NEXT: ret float [[CALL_I]]
//
+// NCRDIV-LABEL: @test___sinf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: ret float [[CALL_I]]
+//
// AMDGCNSPIRV-LABEL: @test___sinf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -6548,6 +7929,14 @@ extern "C" __device__ float test___sinf(float x) {
// APPROX-NEXT: [[MUL_I:%.*]] = fmul contract float [[CALL_I3_I]], [[TMP0]]
// APPROX-NEXT: ret float [[MUL_I]]
//
+// NCRDIV-LABEL: @test___tanf(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[CALL_I3_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR14]]
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rcp.f32(float [[CALL_I_I]])
+// NCRDIV-NEXT: [[MUL_I:%.*]] = fmul contract float [[CALL_I3_I]], [[TMP0]]
+// NCRDIV-NEXT: ret float [[MUL_I]]
+//
// AMDGCNSPIRV-LABEL: @test___tanf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I3_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
@@ -6575,6 +7964,11 @@ extern "C" __device__ float test___tanf(float x) {
// APPROX-NEXT: [[ADD_I:%.*]] = fadd contract double [[X:%.*]], [[Y:%.*]]
// APPROX-NEXT: ret double [[ADD_I]]
//
+// NCRDIV-LABEL: @test___dadd_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[ADD_I:%.*]] = fadd contract double [[X:%.*]], [[Y:%.*]]
+// NCRDIV-NEXT: ret double [[ADD_I]]
+//
// AMDGCNSPIRV-LABEL: @test___dadd_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[ADD_I:%.*]] = fadd contract double [[X:%.*]], [[Y:%.*]]
@@ -6599,6 +7993,11 @@ extern "C" __device__ double test___dadd_rn(double x, double y) {
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract double [[X:%.*]], [[Y:%.*]]
// APPROX-NEXT: ret double [[DIV_I]]
//
+// NCRDIV-LABEL: @test___ddiv_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[DIV_I:%.*]] = fdiv contract double [[X:%.*]], [[Y:%.*]]
+// NCRDIV-NEXT: ret double [[DIV_I]]
+//
// AMDGCNSPIRV-LABEL: @test___ddiv_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[DIV_I:%.*]] = fdiv contract double [[X:%.*]], [[Y:%.*]]
@@ -6623,6 +8022,11 @@ extern "C" __device__ double test___ddiv_rn(double x, double y) {
// APPROX-NEXT: [[MUL_I:%.*]] = fmul contract double [[X:%.*]], [[Y:%.*]]
// APPROX-NEXT: ret double [[MUL_I]]
//
+// NCRDIV-LABEL: @test___dmul_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[MUL_I:%.*]] = fmul contract double [[X:%.*]], [[Y:%.*]]
+// NCRDIV-NEXT: ret double [[MUL_I]]
+//
// AMDGCNSPIRV-LABEL: @test___dmul_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[MUL_I:%.*]] = fmul contract double [[X:%.*]], [[Y:%.*]]
@@ -6647,6 +8051,11 @@ extern "C" __device__ double test___dmul_rn(double x, double y) {
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract double 1.000000e+00, [[X:%.*]]
// APPROX-NEXT: ret double [[DIV_I]]
//
+// NCRDIV-LABEL: @test___drcp_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[DIV_I:%.*]] = fdiv contract double 1.000000e+00, [[X:%.*]]
+// NCRDIV-NEXT: ret double [[DIV_I]]
+//
// AMDGCNSPIRV-LABEL: @test___drcp_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[DIV_I:%.*]] = fdiv contract double 1.000000e+00, [[X:%.*]]
@@ -6671,6 +8080,11 @@ extern "C" __device__ double test___drcp_rn(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[X:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test___dsqrt_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[X:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test___dsqrt_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.sqrt.f64(double [[X:%.*]])
@@ -6695,6 +8109,11 @@ extern "C" __device__ double test___dsqrt_rn(double x) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test__fma_rn(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test__fma_rn(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
@@ -6719,6 +8138,11 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_float_min(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_float_min(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
@@ -6743,6 +8167,11 @@ extern "C" __device__ float test_float_min(float x, float y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
+// NCRDIV-LABEL: @test_float_max(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// NCRDIV-NEXT: ret float [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_float_max(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
@@ -6767,6 +8196,11 @@ extern "C" __device__ float test_float_max(float x, float y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_double_min(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_double_min(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
@@ -6791,6 +8225,11 @@ extern "C" __device__ double test_double_min(double x, double y) {
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
+// NCRDIV-LABEL: @test_double_max(
+// NCRDIV-NEXT: entry:
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// NCRDIV-NEXT: ret double [[TMP0]]
+//
// AMDGCNSPIRV-LABEL: @test_double_max(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
More information about the cfe-commits
mailing list