[clang] 8161656 - HIP: Directly call fabs builtins
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Mon Dec 19 11:20:27 PST 2022
Author: Matt Arsenault
Date: 2022-12-19T14:20:20-05:00
New Revision: 81616561c1b376af85365c9eaf94d49ad184c623
URL: https://github.com/llvm/llvm-project/commit/81616561c1b376af85365c9eaf94d49ad184c623
DIFF: https://github.com/llvm/llvm-project/commit/81616561c1b376af85365c9eaf94d49ad184c623.diff
LOG: HIP: Directly call fabs builtins
All the attribute numbering changes are due to the ocml function being
assumed convergent, and the direct call is not.
This begins a series of patches to avoid using the thinnest
ocml wrappers around llvm intrinsics. These wrapper functions
aren't buying anything and add complexity. As we do not propagate
fast math flags into the library on linking, and cannot generally
safely do so, we're losing information by calling into them. Directly
call the builtins so we get the fast math flags set appropriate for the
translation unit naturally.
Added:
Modified:
clang/lib/Headers/__clang_hip_math.h
clang/test/Headers/__clang_hip_math.hip
clang/test/Headers/amdgcn_openmp_device_math.c
Removed:
################################################################################
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index ef7e087b832ca..186e8212d8ee6 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -233,7 +233,7 @@ __DEVICE__
float expm1f(float __x) { return __ocml_expm1_f32(__x); }
__DEVICE__
-float fabsf(float __x) { return __ocml_fabs_f32(__x); }
+float fabsf(float __x) { return __builtin_fabsf(__x); }
__DEVICE__
float fdimf(float __x, float __y) { return __ocml_fdim_f32(__x, __y); }
@@ -792,7 +792,7 @@ __DEVICE__
double expm1(double __x) { return __ocml_expm1_f64(__x); }
__DEVICE__
-double fabs(double __x) { return __ocml_fabs_f64(__x); }
+double fabs(double __x) { return __builtin_fabs(__x); }
__DEVICE__
double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); }
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index c98c1527b7840..0ad7b2b7d19bd 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -163,12 +163,12 @@ extern "C" __device__ long long test_llabs(long x) {
// DEFAULT-LABEL: @test_acosf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_acosf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_acosf(float x) {
@@ -177,12 +177,12 @@ extern "C" __device__ float test_acosf(float x) {
// DEFAULT-LABEL: @test_acos(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_acos(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_acos(double x) {
@@ -191,12 +191,12 @@ extern "C" __device__ double test_acos(double x) {
// DEFAULT-LABEL: @test_acoshf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_acoshf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_acoshf(float x) {
@@ -205,12 +205,12 @@ extern "C" __device__ float test_acoshf(float x) {
// DEFAULT-LABEL: @test_acosh(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_acosh(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_acosh(double x) {
@@ -219,12 +219,12 @@ extern "C" __device__ double test_acosh(double x) {
// DEFAULT-LABEL: @test_asinf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_asinf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_asinf(float x) {
@@ -233,12 +233,12 @@ extern "C" __device__ float test_asinf(float x) {
// DEFAULT-LABEL: @test_asin(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_asin(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_asin(double x) {
@@ -248,12 +248,12 @@ extern "C" __device__ double test_asin(double x) {
// DEFAULT-LABEL: @test_asinhf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_asinhf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_asinhf(float x) {
@@ -262,12 +262,12 @@ extern "C" __device__ float test_asinhf(float x) {
// DEFAULT-LABEL: @test_asinh(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_asinh(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_asinh(double x) {
@@ -276,12 +276,12 @@ extern "C" __device__ double test_asinh(double x) {
// DEFAULT-LABEL: @test_atan2f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_atan2f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_atan2f(float x, float y) {
@@ -290,12 +290,12 @@ extern "C" __device__ float test_atan2f(float x, float y) {
// DEFAULT-LABEL: @test_atan2(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_atan2(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_atan2(double x, double y) {
@@ -304,12 +304,12 @@ extern "C" __device__ double test_atan2(double x, double y) {
// DEFAULT-LABEL: @test_atanf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_atanf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_atanf(float x) {
@@ -318,12 +318,12 @@ extern "C" __device__ float test_atanf(float x) {
// DEFAULT-LABEL: @test_atan(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_atan(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_atan(double x) {
@@ -332,12 +332,12 @@ extern "C" __device__ double test_atan(double x) {
// DEFAULT-LABEL: @test_atanhf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_atanhf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_atanhf(float x) {
@@ -346,12 +346,12 @@ extern "C" __device__ float test_atanhf(float x) {
// DEFAULT-LABEL: @test_atanh(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_atanh(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_atanh(double x) {
@@ -360,12 +360,12 @@ extern "C" __device__ double test_atanh(double x) {
// DEFAULT-LABEL: @test_cbrtf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_cbrtf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_cbrtf(float x) {
@@ -374,12 +374,12 @@ extern "C" __device__ float test_cbrtf(float x) {
// DEFAULT-LABEL: @test_cbrt(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_cbrt(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_cbrt(double x) {
@@ -388,12 +388,12 @@ extern "C" __device__ double test_cbrt(double x) {
// DEFAULT-LABEL: @test_ceilf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_ceilf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_ceilf(float x) {
@@ -402,12 +402,12 @@ extern "C" __device__ float test_ceilf(float x) {
// DEFAULT-LABEL: @test_ceil(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_ceil(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_ceil(double x) {
@@ -416,12 +416,12 @@ extern "C" __device__ double test_ceil(double x) {
// DEFAULT-LABEL: @test_copysignf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_copysignf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_copysignf(float x, float y) {
@@ -430,12 +430,12 @@ extern "C" __device__ float test_copysignf(float x, float y) {
// DEFAULT-LABEL: @test_copysign(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_copysign(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_copysign(double x, double y) {
@@ -444,12 +444,12 @@ extern "C" __device__ double test_copysign(double x, double y) {
// DEFAULT-LABEL: @test_cosf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR17:[0-9]+]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_cosf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR17:[0-9]+]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_cosf(float x) {
@@ -458,12 +458,12 @@ extern "C" __device__ float test_cosf(float x) {
// DEFAULT-LABEL: @test_cos(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_cos(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_cos(double x) {
@@ -472,12 +472,12 @@ extern "C" __device__ double test_cos(double x) {
// DEFAULT-LABEL: @test_coshf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_coshf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_coshf(float x) {
@@ -486,12 +486,12 @@ extern "C" __device__ float test_coshf(float x) {
// DEFAULT-LABEL: @test_cosh(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_cosh(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_cosh(double x) {
@@ -500,12 +500,12 @@ extern "C" __device__ double test_cosh(double x) {
// DEFAULT-LABEL: @test_cospif(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_cospif(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_cospif(float x) {
@@ -514,12 +514,12 @@ extern "C" __device__ float test_cospif(float x) {
// DEFAULT-LABEL: @test_cospi(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_cospi(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_cospi(double x) {
@@ -528,12 +528,12 @@ extern "C" __device__ double test_cospi(double x) {
// DEFAULT-LABEL: @test_cyl_bessel_i0f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_cyl_bessel_i0f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_cyl_bessel_i0f(float x) {
@@ -542,12 +542,12 @@ extern "C" __device__ float test_cyl_bessel_i0f(float x) {
// DEFAULT-LABEL: @test_cyl_bessel_i0(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_cyl_bessel_i0(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_cyl_bessel_i0(double x) {
@@ -556,12 +556,12 @@ extern "C" __device__ double test_cyl_bessel_i0(double x) {
// DEFAULT-LABEL: @test_cyl_bessel_i1f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_cyl_bessel_i1f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_cyl_bessel_i1f(float x) {
@@ -570,12 +570,12 @@ extern "C" __device__ float test_cyl_bessel_i1f(float x) {
// DEFAULT-LABEL: @test_cyl_bessel_i1(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_cyl_bessel_i1(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_cyl_bessel_i1(double x) {
@@ -584,12 +584,12 @@ extern "C" __device__ double test_cyl_bessel_i1(double x) {
// DEFAULT-LABEL: @test_erfcf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_erfcf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_erfcf(float x) {
@@ -598,12 +598,12 @@ extern "C" __device__ float test_erfcf(float x) {
// DEFAULT-LABEL: @test_erfc(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_erfc(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_erfc(double x) {
@@ -612,12 +612,12 @@ extern "C" __device__ double test_erfc(double x) {
// DEFAULT-LABEL: @test_erfinvf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_erfinvf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_erfinvf(float x) {
@@ -626,12 +626,12 @@ extern "C" __device__ float test_erfinvf(float x) {
// DEFAULT-LABEL: @test_erfinv(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_erfinv(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_erfinv(double x) {
@@ -640,12 +640,12 @@ extern "C" __device__ double test_erfinv(double x) {
// DEFAULT-LABEL: @test_exp10f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_exp10f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_exp10f(float x) {
@@ -654,12 +654,12 @@ extern "C" __device__ float test_exp10f(float x) {
// DEFAULT-LABEL: @test_exp10(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_exp10(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_exp10(double x) {
@@ -668,12 +668,12 @@ extern "C" __device__ double test_exp10(double x) {
// DEFAULT-LABEL: @test_exp2f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_exp2f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_exp2f(float x) {
@@ -682,12 +682,12 @@ extern "C" __device__ float test_exp2f(float x) {
// DEFAULT-LABEL: @test_exp2(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_exp2(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_exp2(double x) {
@@ -696,12 +696,12 @@ extern "C" __device__ double test_exp2(double x) {
// DEFAULT-LABEL: @test_expf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_expf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_expf(float x) {
@@ -710,12 +710,12 @@ extern "C" __device__ float test_expf(float x) {
// DEFAULT-LABEL: @test_exp(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_exp(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_exp(double x) {
@@ -724,12 +724,12 @@ extern "C" __device__ double test_exp(double x) {
// DEFAULT-LABEL: @test_expm1f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_expm1f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_expm1f(float x) {
@@ -738,12 +738,12 @@ extern "C" __device__ float test_expm1f(float x) {
// DEFAULT-LABEL: @test_expm1(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_expm1(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_expm1(double x) {
@@ -752,13 +752,13 @@ extern "C" __device__ double test_expm1(double x) {
// DEFAULT-LABEL: @test_fabsf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fabs_f32(float noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret float [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fabs.f32(float [[X:%.*]])
+// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_fabsf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fabs_f32(float noundef [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret float [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fabs.f32(float [[X:%.*]])
+// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_fabsf(float x) {
return fabsf(x);
@@ -766,13 +766,13 @@ extern "C" __device__ float test_fabsf(float x) {
// DEFAULT-LABEL: @test_fabs(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fabs_f64(double noundef [[X:%.*]]) #[[ATTR14]]
-// DEFAULT-NEXT: ret double [[CALL_I]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fabs.f64(double [[X:%.*]])
+// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_fabs(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fabs_f64(double noundef [[X:%.*]]) #[[ATTR14]]
-// FINITEONLY-NEXT: ret double [[CALL_I]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fabs.f64(double [[X:%.*]])
+// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_fabs(double x) {
return fabs(x);
@@ -780,12 +780,12 @@ extern "C" __device__ double test_fabs(double x) {
// DEFAULT-LABEL: @test_fdimf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_fdimf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_fdimf(float x, float y) {
@@ -794,12 +794,12 @@ extern "C" __device__ float test_fdimf(float x, float y) {
// DEFAULT-LABEL: @test_fdim(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_fdim(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_fdim(double x, double y) {
@@ -822,12 +822,12 @@ extern "C" __device__ float test_fdividef(float x, float y) {
// DEFAULT-LABEL: @test_floorf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_floorf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_floorf(float x) {
@@ -836,12 +836,12 @@ extern "C" __device__ float test_floorf(float x) {
// DEFAULT-LABEL: @test_floor(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_floor(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_floor(double x) {
@@ -850,12 +850,12 @@ extern "C" __device__ double test_floor(double x) {
// DEFAULT-LABEL: @test_fmaf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_fmaf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_fmaf(float x, float y, float z) {
@@ -864,12 +864,12 @@ extern "C" __device__ float test_fmaf(float x, float y, float z) {
// DEFAULT-LABEL: @test_fma(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_fma(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_fma(double x, double y, double z) {
@@ -878,12 +878,12 @@ extern "C" __device__ double test_fma(double x, double y, double z) {
// DEFAULT-LABEL: @test_fma_rn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_fma_rn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_fma_rn(double x, double y, double z) {
@@ -892,12 +892,12 @@ extern "C" __device__ double test_fma_rn(double x, double y, double z) {
// DEFAULT-LABEL: @test_fmaxf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_fmaxf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_fmaxf(float x, float y) {
@@ -906,12 +906,12 @@ extern "C" __device__ float test_fmaxf(float x, float y) {
// DEFAULT-LABEL: @test_fmax(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_fmax(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_fmax(double x, double y) {
@@ -920,12 +920,12 @@ extern "C" __device__ double test_fmax(double x, double y) {
// DEFAULT-LABEL: @test_fminf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_fminf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_fminf(float x, float y) {
@@ -934,12 +934,12 @@ extern "C" __device__ float test_fminf(float x, float y) {
// DEFAULT-LABEL: @test_fmin(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_fmin(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_fmin(double x, double y) {
@@ -948,12 +948,12 @@ extern "C" __device__ double test_fmin(double x, double y) {
// DEFAULT-LABEL: @test_fmodf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_fmodf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_fmodf(float x, float y) {
@@ -962,12 +962,12 @@ extern "C" __device__ float test_fmodf(float x, float y) {
// DEFAULT-LABEL: @test_fmod(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_fmod(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_fmod(double x, double y) {
@@ -977,21 +977,21 @@ extern "C" __device__ double test_fmod(double x, double y) {
// DEFAULT-LABEL: @test_frexpf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17:[0-9]+]]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_frexp_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18:[0-9]+]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_frexp_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11:![0-9]+]]
// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_frexpf(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17:[0-9]+]]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_frexp_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18:[0-9]+]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_frexp_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11:![0-9]+]]
// FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_frexpf(float x, int* y) {
@@ -1001,21 +1001,21 @@ extern "C" __device__ float test_frexpf(float x, int* y) {
// DEFAULT-LABEL: @test_frexp(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_frexp_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_frexp_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_frexp(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_frexp_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_frexp_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
// FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_frexp(double x, int* y) {
@@ -1024,12 +1024,12 @@ extern "C" __device__ double test_frexp(double x, int* y) {
// DEFAULT-LABEL: @test_hypotf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_hypotf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_hypotf(float x, float y) {
@@ -1038,12 +1038,12 @@ extern "C" __device__ float test_hypotf(float x, float y) {
// DEFAULT-LABEL: @test_hypot(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_hypot(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_hypot(double x, double y) {
@@ -1052,7 +1052,7 @@ extern "C" __device__ double test_hypot(double x, double y) {
// CHECK-LABEL: @test_ilogbf(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
+// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]]
// CHECK-NEXT: ret i32 [[CALL_I]]
//
extern "C" __device__ int test_ilogbf(float x) {
@@ -1061,7 +1061,7 @@ extern "C" __device__ int test_ilogbf(float x) {
// CHECK-LABEL: @test_ilogb(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// CHECK-NEXT: ret i32 [[CALL_I]]
//
extern "C" __device__ int test_ilogb(double x) {
@@ -1070,7 +1070,7 @@ extern "C" __device__ int test_ilogb(double x) {
// CHECK-LABEL: @test___finitef(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// CHECK-NEXT: ret i32 [[CONV]]
@@ -1081,7 +1081,7 @@ extern "C" __device__ BOOL_TYPE test___finitef(float x) {
// CHECK-LABEL: @test___finite(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// CHECK-NEXT: ret i32 [[CONV]]
@@ -1092,7 +1092,7 @@ extern "C" __device__ BOOL_TYPE test___finite(double x) {
// CHECK-LABEL: @test___isinff(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// CHECK-NEXT: ret i32 [[CONV]]
@@ -1103,7 +1103,7 @@ extern "C" __device__ BOOL_TYPE test___isinff(float x) {
// CHECK-LABEL: @test___isinf(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// CHECK-NEXT: ret i32 [[CONV]]
@@ -1114,7 +1114,7 @@ extern "C" __device__ BOOL_TYPE test___isinf(double x) {
// CHECK-LABEL: @test___isnanf(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// CHECK-NEXT: ret i32 [[CONV]]
@@ -1125,7 +1125,7 @@ extern "C" __device__ BOOL_TYPE test___isnanf(float x) {
// CHECK-LABEL: @test___isnan(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// CHECK-NEXT: ret i32 [[CONV]]
@@ -1136,12 +1136,12 @@ extern "C" __device__ BOOL_TYPE test___isnan(double x) {
// DEFAULT-LABEL: @test_j0f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_j0f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_j0f(float x) {
@@ -1150,12 +1150,12 @@ extern "C" __device__ float test_j0f(float x) {
// DEFAULT-LABEL: @test_j0(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_j0(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_j0(double x) {
@@ -1164,12 +1164,12 @@ extern "C" __device__ double test_j0(double x) {
// DEFAULT-LABEL: @test_j1f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_j1f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_j1f(float x) {
@@ -1178,12 +1178,12 @@ extern "C" __device__ float test_j1f(float x) {
// DEFAULT-LABEL: @test_j1(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_j1(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_j1(double x) {
@@ -1197,14 +1197,14 @@ extern "C" __device__ double test_j1(double x) {
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR17]]
// DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR17]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR17]]
// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
// DEFAULT: for.body.i:
@@ -1230,14 +1230,14 @@ extern "C" __device__ double test_j1(double x) {
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR17]]
// FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR17]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
// FINITEONLY: for.body.i:
@@ -1267,14 +1267,14 @@ extern "C" __device__ float test_jnf(int x, float y) {
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR17]]
// DEFAULT-NEXT: br label [[_ZL2JNID_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR17]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR17]]
// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
// DEFAULT: for.body.i:
@@ -1300,14 +1300,14 @@ extern "C" __device__ float test_jnf(int x, float y) {
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR17]]
// FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR17]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
// FINITEONLY: for.body.i:
@@ -1332,12 +1332,12 @@ extern "C" __device__ double test_jn(int x, double y) {
// DEFAULT-LABEL: @test_ldexpf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_ldexpf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_ldexpf(float x, int y) {
@@ -1346,12 +1346,12 @@ extern "C" __device__ float test_ldexpf(float x, int y) {
// DEFAULT-LABEL: @test_ldexp(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_ldexp(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_ldexp(double x, int y) {
@@ -1360,12 +1360,12 @@ extern "C" __device__ double test_ldexp(double x, int y) {
// DEFAULT-LABEL: @test_lgammaf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_lgammaf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_lgammaf(float x) {
@@ -1374,12 +1374,12 @@ extern "C" __device__ float test_lgammaf(float x) {
// DEFAULT-LABEL: @test_lgamma(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_lgamma(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_lgamma(double x) {
@@ -1388,13 +1388,13 @@ extern "C" __device__ double test_lgamma(double x) {
// DEFAULT-LABEL: @test_llrintf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_llrintf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -1404,13 +1404,13 @@ extern "C" __device__ long long int test_llrintf(float x) {
// DEFAULT-LABEL: @test_llrint(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_llrint(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -1420,13 +1420,13 @@ extern "C" __device__ long long int test_llrint(double x) {
// DEFAULT-LABEL: @test_llroundf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_llroundf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -1436,13 +1436,13 @@ extern "C" __device__ long long int test_llroundf(float x) {
// DEFAULT-LABEL: @test_llround(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_llround(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -1452,12 +1452,12 @@ extern "C" __device__ long long int test_llround(double x) {
// DEFAULT-LABEL: @test_log10f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_log10f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_log10f(float x) {
@@ -1466,12 +1466,12 @@ extern "C" __device__ float test_log10f(float x) {
// DEFAULT-LABEL: @test_log10(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_log10(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_log10(double x) {
@@ -1480,12 +1480,12 @@ extern "C" __device__ double test_log10(double x) {
// DEFAULT-LABEL: @test_log1pf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_log1pf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_log1pf(float x) {
@@ -1494,12 +1494,12 @@ extern "C" __device__ float test_log1pf(float x) {
// DEFAULT-LABEL: @test_log1p(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_log1p(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_log1p(double x) {
@@ -1508,12 +1508,12 @@ extern "C" __device__ double test_log1p(double x) {
// DEFAULT-LABEL: @test_log2f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_log2f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_log2f(float x) {
@@ -1522,12 +1522,12 @@ extern "C" __device__ float test_log2f(float x) {
// DEFAULT-LABEL: @test_log2(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_log2(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_log2(double x) {
@@ -1536,12 +1536,12 @@ extern "C" __device__ double test_log2(double x) {
// DEFAULT-LABEL: @test_logbf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_logbf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_logbf(float x) {
@@ -1550,12 +1550,12 @@ extern "C" __device__ float test_logbf(float x) {
// DEFAULT-LABEL: @test_logb(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_logb(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_logb(double x) {
@@ -1564,13 +1564,13 @@ extern "C" __device__ double test_logb(double x) {
// DEFAULT-LABEL: @test_lrintf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_lrintf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -1580,13 +1580,13 @@ extern "C" __device__ long int test_lrintf(float x) {
// DEFAULT-LABEL: @test_lrint(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_lrint(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -1596,13 +1596,13 @@ extern "C" __device__ long int test_lrint(double x) {
// DEFAULT-LABEL: @test_lroundf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_lroundf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -1612,13 +1612,13 @@ extern "C" __device__ long int test_lroundf(float x) {
// DEFAULT-LABEL: @test_lround(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]]
//
// FINITEONLY-LABEL: @test_lround(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -1629,21 +1629,21 @@ extern "C" __device__ long int test_lround(double x) {
// DEFAULT-LABEL: @test_modff(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15:![0-9]+]]
// DEFAULT-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_modff(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15:![0-9]+]]
// FINITEONLY-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_modff(float x, float* y) {
@@ -1653,21 +1653,21 @@ extern "C" __device__ float test_modff(float x, float* y) {
// DEFAULT-LABEL: @test_modf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17:![0-9]+]]
// DEFAULT-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_modf(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17:![0-9]+]]
// FINITEONLY-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_modf(double x, double* y) {
@@ -1724,12 +1724,12 @@ extern "C" __device__ double test_nan_fill() {
// DEFAULT-LABEL: @test_nearbyintf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_nearbyintf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_nearbyintf(float x) {
@@ -1738,12 +1738,12 @@ extern "C" __device__ float test_nearbyintf(float x) {
// DEFAULT-LABEL: @test_nearbyint(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_nearbyint(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_nearbyint(double x) {
@@ -1752,12 +1752,12 @@ extern "C" __device__ double test_nearbyint(double x) {
// DEFAULT-LABEL: @test_nextafterf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_nextafterf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_nextafterf(float x, float y) {
@@ -1766,12 +1766,12 @@ extern "C" __device__ float test_nextafterf(float x, float y) {
// DEFAULT-LABEL: @test_nextafter(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_nextafter(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_nextafter(double x, double y) {
@@ -1780,12 +1780,12 @@ extern "C" __device__ double test_nextafter(double x, double y) {
// DEFAULT-LABEL: @test_norm3df(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_norm3df(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_norm3df(float x, float y, float z) {
@@ -1794,12 +1794,12 @@ extern "C" __device__ float test_norm3df(float x, float y, float z) {
// DEFAULT-LABEL: @test_norm3d(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_norm3d(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_norm3d(double x, double y, double z) {
@@ -1808,12 +1808,12 @@ extern "C" __device__ double test_norm3d(double x, double y, double z) {
// DEFAULT-LABEL: @test_norm4df(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_norm4df(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_norm4df(float x, float y, float z, float w) {
@@ -1822,12 +1822,12 @@ extern "C" __device__ float test_norm4df(float x, float y, float z, float w) {
// DEFAULT-LABEL: @test_norm4d(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_norm4d(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_norm4d(double x, double y, double z, double w) {
@@ -1836,12 +1836,12 @@ extern "C" __device__ double test_norm4d(double x, double y, double z, double w)
// DEFAULT-LABEL: @test_normcdff(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_normcdff(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_normcdff(float x) {
@@ -1850,12 +1850,12 @@ extern "C" __device__ float test_normcdff(float x) {
// DEFAULT-LABEL: @test_normcdf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_normcdf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_normcdf(double x) {
@@ -1864,12 +1864,12 @@ extern "C" __device__ double test_normcdf(double x) {
// DEFAULT-LABEL: @test_normcdfinvf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_normcdfinvf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_normcdfinvf(float x) {
@@ -1878,12 +1878,12 @@ extern "C" __device__ float test_normcdfinvf(float x) {
// DEFAULT-LABEL: @test_normcdfinv(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_normcdfinv(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_normcdfinv(double x) {
@@ -1907,7 +1907,7 @@ extern "C" __device__ double test_normcdfinv(double x) {
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP19:![0-9]+]]
// DEFAULT: _ZL5normfiPKf.exit:
// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_normf(
@@ -1927,7 +1927,7 @@ extern "C" __device__ double test_normcdfinv(double x) {
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP19:![0-9]+]]
// FINITEONLY: _ZL5normfiPKf.exit:
// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_normf(int x, const float *y) {
@@ -1951,7 +1951,7 @@ extern "C" __device__ float test_normf(int x, const float *y) {
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
// DEFAULT: _ZL4normiPKd.exit:
// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_norm(
@@ -1971,7 +1971,7 @@ extern "C" __device__ float test_normf(int x, const float *y) {
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
// FINITEONLY: _ZL4normiPKd.exit:
// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_norm(int x, const double *y) {
@@ -1980,12 +1980,12 @@ extern "C" __device__ double test_norm(int x, const double *y) {
// DEFAULT-LABEL: @test_powf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_powf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_powf(float x, float y) {
@@ -1994,12 +1994,12 @@ extern "C" __device__ float test_powf(float x, float y) {
// DEFAULT-LABEL: @test_pow(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_pow(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_pow(double x, double y) {
@@ -2008,12 +2008,12 @@ extern "C" __device__ double test_pow(double x, double y) {
// DEFAULT-LABEL: @test_powif(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_powif(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_powif(float x, int y) {
@@ -2022,12 +2022,12 @@ extern "C" __device__ float test_powif(float x, int y) {
// DEFAULT-LABEL: @test_powi(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_powi(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_powi(double x, int y) {
@@ -2036,12 +2036,12 @@ extern "C" __device__ double test_powi(double x, int y) {
// DEFAULT-LABEL: @test_rcbrtf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rcbrtf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_rcbrtf(float x) {
@@ -2050,12 +2050,12 @@ extern "C" __device__ float test_rcbrtf(float x) {
// DEFAULT-LABEL: @test_rcbrt(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rcbrt(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_rcbrt(double x) {
@@ -2064,12 +2064,12 @@ extern "C" __device__ double test_rcbrt(double x) {
// DEFAULT-LABEL: @test_remainderf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_remainderf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_remainderf(float x, float y) {
@@ -2078,12 +2078,12 @@ extern "C" __device__ float test_remainderf(float x, float y) {
// DEFAULT-LABEL: @test_remainder(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_remainder(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_remainder(double x, double y) {
@@ -2093,21 +2093,21 @@ extern "C" __device__ double test_remainder(double x, double y) {
// DEFAULT-LABEL: @test_remquof(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_remquof(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
// FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_remquof(float x, float y, int* z) {
@@ -2117,21 +2117,21 @@ extern "C" __device__ float test_remquof(float x, float y, int* z) {
// DEFAULT-LABEL: @test_remquo(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_remquo(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
// FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_remquo(double x, double y, int* z) {
@@ -2140,12 +2140,12 @@ extern "C" __device__ double test_remquo(double x, double y, int* z) {
// DEFAULT-LABEL: @test_rhypotf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rhypotf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_rhypotf(float x, float y) {
@@ -2154,12 +2154,12 @@ extern "C" __device__ float test_rhypotf(float x, float y) {
// DEFAULT-LABEL: @test_rhypot(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rhypot(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_rhypot(double x, double y) {
@@ -2168,12 +2168,12 @@ extern "C" __device__ double test_rhypot(double x, double y) {
// DEFAULT-LABEL: @test_rintf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rintf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_rintf(float x) {
@@ -2182,12 +2182,12 @@ extern "C" __device__ float test_rintf(float x) {
// DEFAULT-LABEL: @test_rint(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rint(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_rint(double x) {
@@ -2211,7 +2211,7 @@ extern "C" __device__ double test_rint(double x) {
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
// DEFAULT: _ZL6rnormfiPKf.exit:
// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnormf(
@@ -2231,7 +2231,7 @@ extern "C" __device__ double test_rint(double x) {
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
// FINITEONLY: _ZL6rnormfiPKf.exit:
// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_rnormf(int x, const float* y) {
@@ -2255,7 +2255,7 @@ extern "C" __device__ float test_rnormf(int x, const float* y) {
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
// DEFAULT: _ZL5rnormiPKd.exit:
// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnorm(
@@ -2275,7 +2275,7 @@ extern "C" __device__ float test_rnormf(int x, const float* y) {
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
// FINITEONLY: _ZL5rnormiPKd.exit:
// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_rnorm(int x, const double* y) {
@@ -2284,12 +2284,12 @@ extern "C" __device__ double test_rnorm(int x, const double* y) {
// DEFAULT-LABEL: @test_rnorm3df(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnorm3df(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_rnorm3df(float x, float y, float z) {
@@ -2298,12 +2298,12 @@ extern "C" __device__ float test_rnorm3df(float x, float y, float z) {
// DEFAULT-LABEL: @test_rnorm3d(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnorm3d(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_rnorm3d(double x, double y, double z) {
@@ -2312,12 +2312,12 @@ extern "C" __device__ double test_rnorm3d(double x, double y, double z) {
// DEFAULT-LABEL: @test_rnorm4df(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnorm4df(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_rnorm4df(float x, float y, float z, float w) {
@@ -2326,12 +2326,12 @@ extern "C" __device__ float test_rnorm4df(float x, float y, float z, float w) {
// DEFAULT-LABEL: @test_rnorm4d(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnorm4d(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w) {
@@ -2340,12 +2340,12 @@ extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w
// DEFAULT-LABEL: @test_roundf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_roundf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_roundf(float x) {
@@ -2354,12 +2354,12 @@ extern "C" __device__ float test_roundf(float x) {
// DEFAULT-LABEL: @test_round(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_round(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_round(double x) {
@@ -2368,12 +2368,12 @@ extern "C" __device__ double test_round(double x) {
// DEFAULT-LABEL: @test_rsqrtf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rsqrtf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_rsqrtf(float x) {
@@ -2382,12 +2382,12 @@ extern "C" __device__ float test_rsqrtf(float x) {
// DEFAULT-LABEL: @test_rsqrt(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rsqrt(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_rsqrt(double x) {
@@ -2400,10 +2400,10 @@ extern "C" __device__ double test_rsqrt(double x) {
// DEFAULT-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
// DEFAULT: cond.true.i:
// DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL8SCALBLNFFL_EXIT:%.*]]
// DEFAULT: cond.false.i:
-// DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL8SCALBLNFFL_EXIT]]
// DEFAULT: _ZL8scalblnffl.exit:
// DEFAULT-NEXT: [[COND_I:%.*]] = phi contract float [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
@@ -2415,10 +2415,10 @@ extern "C" __device__ double test_rsqrt(double x) {
// FINITEONLY-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
// FINITEONLY: cond.true.i:
// FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL8SCALBLNFFL_EXIT:%.*]]
// FINITEONLY: cond.false.i:
-// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL8SCALBLNFFL_EXIT]]
// FINITEONLY: _ZL8scalblnffl.exit:
// FINITEONLY-NEXT: [[COND_I:%.*]] = phi nnan ninf contract float [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
@@ -2434,10 +2434,10 @@ extern "C" __device__ float test_scalblnf(float x, long int y) {
// DEFAULT-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
// DEFAULT: cond.true.i:
// DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL7SCALBLNDL_EXIT:%.*]]
// DEFAULT: cond.false.i:
-// DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL7SCALBLNDL_EXIT]]
// DEFAULT: _ZL7scalblndl.exit:
// DEFAULT-NEXT: [[COND_I:%.*]] = phi contract double [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
@@ -2449,10 +2449,10 @@ extern "C" __device__ float test_scalblnf(float x, long int y) {
// FINITEONLY-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
// FINITEONLY: cond.true.i:
// FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL7SCALBLNDL_EXIT:%.*]]
// FINITEONLY: cond.false.i:
-// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL7SCALBLNDL_EXIT]]
// FINITEONLY: _ZL7scalblndl.exit:
// FINITEONLY-NEXT: [[COND_I:%.*]] = phi nnan ninf contract double [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
@@ -2464,12 +2464,12 @@ extern "C" __device__ double test_scalbln(double x, long int y) {
// DEFAULT-LABEL: @test_scalbnf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_scalbnf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_scalbnf(float x, int y) {
@@ -2478,12 +2478,12 @@ extern "C" __device__ float test_scalbnf(float x, int y) {
// DEFAULT-LABEL: @test_scalbn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_scalbn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_scalbn(double x, int y) {
@@ -2492,7 +2492,7 @@ extern "C" __device__ double test_scalbn(double x, int y) {
// CHECK-LABEL: @test___signbitf(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// CHECK-NEXT: ret i32 [[CONV]]
@@ -2503,7 +2503,7 @@ extern "C" __device__ BOOL_TYPE test___signbitf(float x) {
// CHECK-LABEL: @test___signbit(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// CHECK-NEXT: ret i32 [[CONV]]
@@ -2515,23 +2515,23 @@ extern "C" __device__ BOOL_TYPE test___signbit(double x) {
// DEFAULT-LABEL: @test_sincosf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]]
// DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// DEFAULT-NEXT: ret void
//
// FINITEONLY-LABEL: @test_sincosf(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]]
// FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret void
//
extern "C" __device__ void test_sincosf(float x, float *y, float *z) {
@@ -2541,23 +2541,23 @@ extern "C" __device__ void test_sincosf(float x, float *y, float *z) {
// DEFAULT-LABEL: @test_sincos(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]]
// DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// DEFAULT-NEXT: ret void
//
// FINITEONLY-LABEL: @test_sincos(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]]
// FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret void
//
extern "C" __device__ void test_sincos(double x, double *y, double *z) {
@@ -2567,23 +2567,23 @@ extern "C" __device__ void test_sincos(double x, double *y, double *z) {
// DEFAULT-LABEL: @test_sincospif(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]]
// DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// DEFAULT-NEXT: ret void
//
// FINITEONLY-LABEL: @test_sincospif(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]]
// FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret void
//
extern "C" __device__ void test_sincospif(float x, float *y, float *z) {
@@ -2593,23 +2593,23 @@ extern "C" __device__ void test_sincospif(float x, float *y, float *z) {
// DEFAULT-LABEL: @test_sincospi(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]]
// DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// DEFAULT-NEXT: ret void
//
// FINITEONLY-LABEL: @test_sincospi(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR17]]
// FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]]
// FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
// FINITEONLY-NEXT: ret void
//
extern "C" __device__ void test_sincospi(double x, double *y, double *z) {
@@ -2618,12 +2618,12 @@ extern "C" __device__ void test_sincospi(double x, double *y, double *z) {
// DEFAULT-LABEL: @test_sinf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_sinf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_sinf(float x) {
@@ -2632,12 +2632,12 @@ extern "C" __device__ float test_sinf(float x) {
// DEFAULT-LABEL: @test_sin(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_sin(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_sin(double x) {
@@ -2646,12 +2646,12 @@ extern "C" __device__ double test_sin(double x) {
// DEFAULT-LABEL: @test_sinpif(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_sinpif(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_sinpif(float x) {
@@ -2660,12 +2660,12 @@ extern "C" __device__ float test_sinpif(float x) {
// DEFAULT-LABEL: @test_sinpi(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_sinpi(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_sinpi(double x) {
@@ -2674,12 +2674,12 @@ extern "C" __device__ double test_sinpi(double x) {
// DEFAULT-LABEL: @test_sqrtf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_sqrtf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_sqrtf(float x) {
@@ -2688,12 +2688,12 @@ extern "C" __device__ float test_sqrtf(float x) {
// DEFAULT-LABEL: @test_sqrt(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_sqrt(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_sqrt(double x) {
@@ -2702,12 +2702,12 @@ extern "C" __device__ double test_sqrt(double x) {
// DEFAULT-LABEL: @test_tanf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_tanf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_tanf(float x) {
@@ -2716,12 +2716,12 @@ extern "C" __device__ float test_tanf(float x) {
// DEFAULT-LABEL: @test_tan(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_tan(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_tan(double x) {
@@ -2730,12 +2730,12 @@ extern "C" __device__ double test_tan(double x) {
// DEFAULT-LABEL: @test_tanhf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_tanhf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_tanhf(float x) {
@@ -2744,12 +2744,12 @@ extern "C" __device__ float test_tanhf(float x) {
// DEFAULT-LABEL: @test_tanh(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_tanh(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_tanh(double x) {
@@ -2758,12 +2758,12 @@ extern "C" __device__ double test_tanh(double x) {
// DEFAULT-LABEL: @test_tgammaf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_tgammaf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_tgammaf(float x) {
@@ -2772,12 +2772,12 @@ extern "C" __device__ float test_tgammaf(float x) {
// DEFAULT-LABEL: @test_tgamma(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_tgamma(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_tgamma(double x) {
@@ -2786,12 +2786,12 @@ extern "C" __device__ double test_tgamma(double x) {
// DEFAULT-LABEL: @test_truncf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_truncf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_truncf(float x) {
@@ -2800,12 +2800,12 @@ extern "C" __device__ float test_truncf(float x) {
// DEFAULT-LABEL: @test_trunc(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_trunc(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_trunc(double x) {
@@ -2814,12 +2814,12 @@ extern "C" __device__ double test_trunc(double x) {
// DEFAULT-LABEL: @test_y0f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_y0f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_y0f(float x) {
@@ -2828,12 +2828,12 @@ extern "C" __device__ float test_y0f(float x) {
// DEFAULT-LABEL: @test_y0(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_y0(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_y0(double x) {
@@ -2842,12 +2842,12 @@ extern "C" __device__ double test_y0(double x) {
// DEFAULT-LABEL: @test_y1f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_y1f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_y1f(float x) {
@@ -2856,12 +2856,12 @@ extern "C" __device__ float test_y1f(float x) {
// DEFAULT-LABEL: @test_y1(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_y1(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_y1(double x) {
@@ -2875,14 +2875,14 @@ extern "C" __device__ double test_y1(double x) {
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR17]]
// DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR17]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR17]]
// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
// DEFAULT: for.body.i:
@@ -2908,14 +2908,14 @@ extern "C" __device__ double test_y1(double x) {
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR17]]
// FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR17]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
// FINITEONLY: for.body.i:
@@ -2945,14 +2945,14 @@ extern "C" __device__ float test_ynf(int x, float y) {
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR17]]
// DEFAULT-NEXT: br label [[_ZL2YNID_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR17]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR17]]
// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
// DEFAULT: for.body.i:
@@ -2978,14 +2978,14 @@ extern "C" __device__ float test_ynf(int x, float y) {
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR17]]
// FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR17]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR17]]
// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
// FINITEONLY: for.body.i:
@@ -3010,12 +3010,12 @@ extern "C" __device__ double test_yn(int x, double y) {
// DEFAULT-LABEL: @test___cosf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test___cosf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test___cosf(float x) {
@@ -3024,12 +3024,12 @@ extern "C" __device__ float test___cosf(float x) {
// DEFAULT-LABEL: @test___exp10f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test___exp10f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test___exp10f(float x) {
@@ -3038,12 +3038,12 @@ extern "C" __device__ float test___exp10f(float x) {
// DEFAULT-LABEL: @test___expf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test___expf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test___expf(float x) {
@@ -3080,12 +3080,12 @@ extern "C" __device__ float test___fdividef(float x, float y) {
// DEFAULT-LABEL: @test__fmaf_rn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test__fmaf_rn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
@@ -3122,12 +3122,12 @@ extern "C" __device__ float test___frcp_rn(float x) {
// DEFAULT-LABEL: @test___fsqrt_rn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test___fsqrt_rn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test___fsqrt_rn(float x) {
@@ -3150,12 +3150,12 @@ extern "C" __device__ float test___fsub_rn(float x, float y) {
// DEFAULT-LABEL: @test___log10f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test___log10f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test___log10f(float x) {
@@ -3164,12 +3164,12 @@ extern "C" __device__ float test___log10f(float x) {
// DEFAULT-LABEL: @test___log2f(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test___log2f(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test___log2f(float x) {
@@ -3178,12 +3178,12 @@ extern "C" __device__ float test___log2f(float x) {
// DEFAULT-LABEL: @test___logf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test___logf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test___logf(float x) {
@@ -3192,12 +3192,12 @@ extern "C" __device__ float test___logf(float x) {
// DEFAULT-LABEL: @test___powf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test___powf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test___powf(float x, float y) {
@@ -3226,17 +3226,17 @@ extern "C" __device__ float test___saturatef(float x) {
// DEFAULT-LABEL: @test___sincosf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
-// DEFAULT-NEXT: [[CALL1_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL1_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR17]]
// DEFAULT-NEXT: store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
// DEFAULT-NEXT: ret void
//
// FINITEONLY-LABEL: @test___sincosf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
-// FINITEONLY-NEXT: [[CALL1_I:%.*]] = tail call nnan ninf contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL1_I:%.*]] = tail call nnan ninf contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR17]]
// FINITEONLY-NEXT: store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
// FINITEONLY-NEXT: ret void
//
@@ -3246,12 +3246,12 @@ extern "C" __device__ void test___sincosf(float x, float *y, float *z) {
// DEFAULT-LABEL: @test___sinf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test___sinf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test___sinf(float x) {
@@ -3260,12 +3260,12 @@ extern "C" __device__ float test___sinf(float x) {
// DEFAULT-LABEL: @test___tanf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test___tanf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR17]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test___tanf(float x) {
@@ -3330,12 +3330,12 @@ extern "C" __device__ double test___drcp_rn(double x) {
// DEFAULT-LABEL: @test___dsqrt_rn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test___dsqrt_rn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test___dsqrt_rn(double x) {
@@ -3344,12 +3344,12 @@ extern "C" __device__ double test___dsqrt_rn(double x) {
// DEFAULT-LABEL: @test__fma_rn(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test__fma_rn(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test__fma_rn(double x, double y, double z) {
@@ -3358,12 +3358,12 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) {
// DEFAULT-LABEL: @test_float_min(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I_I]]
//
// FINITEONLY-LABEL: @test_float_min(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I_I]]
//
extern "C" __device__ float test_float_min(float x, float y) {
@@ -3372,12 +3372,12 @@ extern "C" __device__ float test_float_min(float x, float y) {
// DEFAULT-LABEL: @test_float_max(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I_I]]
//
// FINITEONLY-LABEL: @test_float_max(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I_I]]
//
extern "C" __device__ float test_float_max(float x, float y) {
@@ -3386,12 +3386,12 @@ extern "C" __device__ float test_float_max(float x, float y) {
// DEFAULT-LABEL: @test_double_min(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I_I]]
//
// FINITEONLY-LABEL: @test_double_min(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I_I]]
//
extern "C" __device__ double test_double_min(double x, double y) {
@@ -3400,12 +3400,12 @@ extern "C" __device__ double test_double_min(double x, double y) {
// DEFAULT-LABEL: @test_double_max(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I_I]]
//
// FINITEONLY-LABEL: @test_double_max(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I_I]]
//
extern "C" __device__ double test_double_max(double x, double y) {
diff --git a/clang/test/Headers/amdgcn_openmp_device_math.c b/clang/test/Headers/amdgcn_openmp_device_math.c
index cab1e88156f67..1f3eac82c038a 100644
--- a/clang/test/Headers/amdgcn_openmp_device_math.c
+++ b/clang/test/Headers/amdgcn_openmp_device_math.c
@@ -17,7 +17,7 @@ void test_math_f64(double x) {
double l1 = sin(x);
// CHECK: call double @__ocml_cos_f64
double l2 = cos(x);
- // CHECK: call double @__ocml_fabs_f64
+ // CHECK: call double @llvm.fabs.f64
double l3 = fabs(x);
}
}
@@ -32,8 +32,8 @@ void test_math_f32(float x) {
// CHECK-C: call double @__ocml_cos_f64
// CHECK-CPP: call float @__ocml_cos_f32
float l2 = cos(x);
- // CHECK-C: call double @__ocml_fabs_f64
- // CHECK-CPP: call float @__ocml_fabs_f32
+ // CHECK-C: call double @llvm.fabs.f64
+ // CHECK-CPP: call float @llvm.fabs.f32
float l3 = fabs(x);
}
}
@@ -45,7 +45,7 @@ void test_math_f32_suffix(float x) {
float l1 = sinf(x);
// CHECK: call float @__ocml_cos_f32
float l2 = cosf(x);
- // CHECK: call float @__ocml_fabs_f32
+ // CHECK: call float @llvm.fabs.f32
float l3 = fabsf(x);
}
}
More information about the cfe-commits
mailing list