Index: clang/lib/Headers/__clang_hip_math.h =================================================================== --- clang/lib/Headers/__clang_hip_math.h +++ clang/lib/Headers/__clang_hip_math.h @@ -242,7 +242,7 @@ 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); } @@ -801,7 +801,7 @@ 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); } Index: clang/test/Headers/__clang_hip_math.hip =================================================================== --- clang/test/Headers/__clang_hip_math.hip +++ clang/test/Headers/__clang_hip_math.hip @@ -256,12 +256,12 @@ // 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) { @@ -270,12 +270,12 @@ // 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) { @@ -284,12 +284,12 @@ // 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) { @@ -298,12 +298,12 @@ // 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) { @@ -312,12 +312,12 @@ // 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) { @@ -326,12 +326,12 @@ // 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) { @@ -341,12 +341,12 @@ // 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) { @@ -355,12 +355,12 @@ // 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) { @@ -369,12 +369,12 @@ // 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) { @@ -383,12 +383,12 @@ // 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) { @@ -397,12 +397,12 @@ // 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) { @@ -411,12 +411,12 @@ // 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) { @@ -425,12 +425,12 @@ // 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) { @@ -439,12 +439,12 @@ // 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) { @@ -453,12 +453,12 @@ // 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) { @@ -467,12 +467,12 @@ // 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) { @@ -481,12 +481,12 @@ // 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) { @@ -495,12 +495,12 @@ // 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) { @@ -509,12 +509,12 @@ // 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) { @@ -523,12 +523,12 @@ // 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) { @@ -537,12 +537,12 @@ // 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) { @@ -551,12 +551,12 @@ // 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) { @@ -565,12 +565,12 @@ // 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) { @@ -579,12 +579,12 @@ // 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) { @@ -593,12 +593,12 @@ // 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) { @@ -607,12 +607,12 @@ // 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) { @@ -621,12 +621,12 @@ // 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) { @@ -635,12 +635,12 @@ // 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) { @@ -649,12 +649,12 @@ // 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) { @@ -663,12 +663,12 @@ // 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) { @@ -677,12 +677,12 @@ // 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) { @@ -691,12 +691,12 @@ // 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) { @@ -705,12 +705,12 @@ // 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) { @@ -719,12 +719,12 @@ // 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) { @@ -733,12 +733,12 @@ // 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) { @@ -747,12 +747,12 @@ // 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) { @@ -761,12 +761,12 @@ // 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) { @@ -775,12 +775,12 @@ // 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) { @@ -789,12 +789,12 @@ // 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) { @@ -803,12 +803,12 @@ // 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) { @@ -817,12 +817,12 @@ // 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) { @@ -831,12 +831,12 @@ // 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) { @@ -845,13 +845,13 @@ // 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); @@ -859,13 +859,13 @@ // 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); @@ -873,12 +873,12 @@ // 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) { @@ -887,12 +887,12 @@ // 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) { @@ -915,12 +915,12 @@ // 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) { @@ -929,12 +929,12 @@ // 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) { @@ -943,12 +943,12 @@ // 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) { @@ -957,12 +957,12 @@ // 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) { @@ -971,12 +971,12 @@ // 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) { @@ -985,12 +985,12 @@ // 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) { @@ -999,12 +999,12 @@ // 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) { @@ -1013,12 +1013,12 @@ // 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) { @@ -1027,12 +1027,12 @@ // 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) { @@ -1041,12 +1041,12 @@ // 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) { @@ -1055,12 +1055,12 @@ // 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) { @@ -1069,12 +1069,12 @@ // 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) { @@ -1083,12 +1083,12 @@ // 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) { @@ -1097,7 +1097,7 @@ // 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) { @@ -1106,7 +1106,7 @@ // 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) { @@ -1115,7 +1115,7 @@ // 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]] @@ -1126,7 +1126,7 @@ // 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]] @@ -1137,7 +1137,7 @@ // 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]] @@ -1148,7 +1148,7 @@ // 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]] @@ -1159,7 +1159,7 @@ // 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]] @@ -1170,7 +1170,7 @@ // 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]] @@ -1181,12 +1181,12 @@ // 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) { @@ -1195,12 +1195,12 @@ // 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) { @@ -1209,12 +1209,12 @@ // 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) { @@ -1223,12 +1223,12 @@ // 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) { @@ -1242,14 +1242,14 @@ // 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: @@ -1275,14 +1275,14 @@ // 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: @@ -1312,14 +1312,14 @@ // 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: @@ -1345,14 +1345,14 @@ // 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: @@ -1377,12 +1377,12 @@ // 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) { @@ -1391,12 +1391,12 @@ // 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) { @@ -1405,12 +1405,12 @@ // 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) { @@ -1419,12 +1419,12 @@ // 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) { @@ -1433,13 +1433,13 @@ // 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]] // @@ -1449,13 +1449,13 @@ // 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]] // @@ -1465,13 +1465,13 @@ // 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]] // @@ -1481,13 +1481,13 @@ // 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]] // @@ -1497,12 +1497,12 @@ // 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) { @@ -1511,12 +1511,12 @@ // 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) { @@ -1525,12 +1525,12 @@ // 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) { @@ -1539,12 +1539,12 @@ // 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) { @@ -1553,12 +1553,12 @@ // 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) { @@ -1567,12 +1567,12 @@ // 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) { @@ -1581,12 +1581,12 @@ // 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) { @@ -1595,12 +1595,12 @@ // 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) { @@ -1609,13 +1609,13 @@ // 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]] // @@ -1625,13 +1625,13 @@ // 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]] // @@ -1641,13 +1641,13 @@ // 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]] // @@ -1657,13 +1657,13 @@ // 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]] // @@ -1674,21 +1674,21 @@ // 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:[0-9]+]] -// 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:[0-9]+]] +// 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 [[TBAA13:![0-9]+]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] -// 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:[0-9]+]] -// 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:[0-9]+]] +// 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 [[TBAA13:![0-9]+]] // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] -// 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) { @@ -1698,21 +1698,21 @@ // 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 [[TBAA15:![0-9]+]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] -// 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 [[TBAA15:![0-9]+]] // FINITEONLY-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] -// 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) { @@ -1956,12 +1956,12 @@ // 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) { @@ -1970,12 +1970,12 @@ // 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) { @@ -1984,12 +1984,12 @@ // 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) { @@ -1998,12 +1998,12 @@ // 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) { @@ -2012,12 +2012,12 @@ // 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) { @@ -2026,12 +2026,12 @@ // 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) { @@ -2040,12 +2040,12 @@ // 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) { @@ -2054,12 +2054,12 @@ // 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) { @@ -2068,12 +2068,12 @@ // 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) { @@ -2082,12 +2082,12 @@ // 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) { @@ -2096,12 +2096,12 @@ // 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) { @@ -2110,12 +2110,12 @@ // 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) { @@ -2139,7 +2139,7 @@ // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP17:![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( @@ -2159,7 +2159,7 @@ // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP17:![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) { @@ -2183,7 +2183,7 @@ // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP18:![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( @@ -2203,7 +2203,7 @@ // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP18:![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) { @@ -2212,12 +2212,12 @@ // 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) { @@ -2226,12 +2226,12 @@ // 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) { @@ -2240,12 +2240,12 @@ // 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) { @@ -2254,12 +2254,12 @@ // 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) { @@ -2268,12 +2268,12 @@ // 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) { @@ -2282,12 +2282,12 @@ // 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) { @@ -2296,12 +2296,12 @@ // 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) { @@ -2310,12 +2310,12 @@ // 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) { @@ -2325,21 +2325,21 @@ // 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 [[TBAA19:![0-9]+]] // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]] -// 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 [[TBAA19:![0-9]+]] // FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]] -// 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) { @@ -2349,21 +2349,21 @@ // 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 [[TBAA19]] // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]] -// 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 [[TBAA19]] // FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]] -// 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) { @@ -2372,12 +2372,12 @@ // 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) { @@ -2386,12 +2386,12 @@ // 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) { @@ -2400,12 +2400,12 @@ // 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) { @@ -2414,12 +2414,12 @@ // 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) { @@ -2443,7 +2443,7 @@ // 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( @@ -2463,7 +2463,7 @@ // 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) { @@ -2487,7 +2487,7 @@ // 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( @@ -2507,7 +2507,7 @@ // 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) { @@ -2516,12 +2516,12 @@ // 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) { @@ -2530,12 +2530,12 @@ // 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) { @@ -2544,12 +2544,12 @@ // 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) { @@ -2558,12 +2558,12 @@ // 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) { @@ -2572,12 +2572,12 @@ // 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) { @@ -2586,12 +2586,12 @@ // 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) { @@ -2600,12 +2600,12 @@ // 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) { @@ -2614,12 +2614,12 @@ // 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) { @@ -2632,10 +2632,10 @@ // 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]] ] @@ -2647,10 +2647,10 @@ // 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]] ] @@ -2666,10 +2666,10 @@ // 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]] ] @@ -2681,10 +2681,10 @@ // 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]] ] @@ -2696,12 +2696,12 @@ // 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) { @@ -2710,12 +2710,12 @@ // 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) { @@ -2724,7 +2724,7 @@ // 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]] @@ -2735,7 +2735,7 @@ // 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]] @@ -2747,23 +2747,23 @@ // 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 [[TBAA13]] // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] -// 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 [[TBAA13]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] -// 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) { @@ -2773,23 +2773,23 @@ // 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 [[TBAA15]] // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]] -// 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 [[TBAA15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]] // FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]] -// 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) { @@ -2799,23 +2799,23 @@ // 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 [[TBAA13]] // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] -// 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 [[TBAA13]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] -// 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) { @@ -2825,23 +2825,23 @@ // 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 [[TBAA15]] // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]] -// 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 [[TBAA15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]] // FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]] -// 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) { @@ -2850,12 +2850,12 @@ // 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) { @@ -2864,12 +2864,12 @@ // 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) { @@ -2878,12 +2878,12 @@ // 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) { @@ -2892,12 +2892,12 @@ // 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) { @@ -2906,12 +2906,12 @@ // 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) { @@ -2920,12 +2920,12 @@ // 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) { @@ -2934,12 +2934,12 @@ // 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) { @@ -2948,12 +2948,12 @@ // 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) { @@ -2962,12 +2962,12 @@ // 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) { @@ -2976,12 +2976,12 @@ // 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) { @@ -2990,12 +2990,12 @@ // 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) { @@ -3004,12 +3004,12 @@ // 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) { @@ -3018,12 +3018,12 @@ // 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) { @@ -3032,12 +3032,12 @@ // 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) { @@ -3046,12 +3046,12 @@ // 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) { @@ -3060,12 +3060,12 @@ // 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) { @@ -3074,12 +3074,12 @@ // 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) { @@ -3088,12 +3088,12 @@ // 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) { @@ -3107,14 +3107,14 @@ // 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: @@ -3140,14 +3140,14 @@ // 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: @@ -3177,14 +3177,14 @@ // 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: @@ -3210,14 +3210,14 @@ // 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: @@ -3242,12 +3242,12 @@ // 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) { @@ -3256,12 +3256,12 @@ // 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) { @@ -3270,12 +3270,12 @@ // 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) { @@ -3312,12 +3312,12 @@ // 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) { @@ -3354,12 +3354,12 @@ // 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) { @@ -3382,12 +3382,12 @@ // 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) { @@ -3396,12 +3396,12 @@ // 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) { @@ -3410,12 +3410,12 @@ // 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) { @@ -3424,12 +3424,12 @@ // 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) { @@ -3458,17 +3458,17 @@ // 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 [[TBAA13]] -// 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 [[TBAA13]] // 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 [[TBAA13]] -// 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 [[TBAA13]] // FINITEONLY-NEXT: ret void // @@ -3478,12 +3478,12 @@ // 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) { @@ -3492,12 +3492,12 @@ // 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) { @@ -3562,12 +3562,12 @@ // 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) { @@ -3576,12 +3576,12 @@ // 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) { @@ -3590,12 +3590,12 @@ // 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) { @@ -3604,12 +3604,12 @@ // 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) { @@ -3618,12 +3618,12 @@ // 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) { @@ -3632,12 +3632,12 @@ // 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) {