diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h index a4e557e45e5a0..71cda10859ce2 100644 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -257,15 +257,8 @@ float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); } __DEVICE__ float frexpf(float __x, int *__nptr) { - int __tmp; -#ifdef __OPENMP_AMDGCN__ -#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc) -#endif - float __r = - __ocml_frexp_f32(__x, (__attribute__((address_space(5))) int *)&__tmp); - *__nptr = __tmp; - - return __r; + *__nptr = __builtin_amdgcn_frexp_expf(__x); + return __builtin_amdgcn_frexp_mantf(__x); } __DEVICE__ @@ -813,14 +806,8 @@ double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); } __DEVICE__ double frexp(double __x, int *__nptr) { - int __tmp; -#ifdef __OPENMP_AMDGCN__ -#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc) -#endif - double __r = - __ocml_frexp_f64(__x, (__attribute__((address_space(5))) int *)&__tmp); - *__nptr = __tmp; - return __r; + *__nptr = __builtin_amdgcn_frexp_exp(__x); + return __builtin_amdgcn_frexp_mant(__x); } __DEVICE__ diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index 2a60cbaf5f440..d4f87cd4f436a 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -250,12 +250,12 @@ extern "C" __device__ long long test_llabs(long x) { // DEFAULT-LABEL: @test_acosf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR13:[0-9]+]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_acosf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_acos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13:[0-9]+]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_acos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14:[0-9]+]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_acosf(float x) { @@ -264,12 +264,12 @@ extern "C" __device__ float test_acosf(float x) { // DEFAULT-LABEL: @test_acos( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_acos( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_acos_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_acos_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_acos(double x) { @@ -278,12 +278,12 @@ extern "C" __device__ double test_acos(double x) { // DEFAULT-LABEL: @test_acoshf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_acoshf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_acosh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14:[0-9]+]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_acosh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15:[0-9]+]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_acoshf(float x) { @@ -292,12 +292,12 @@ extern "C" __device__ float test_acoshf(float x) { // DEFAULT-LABEL: @test_acosh( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_acosh( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_acosh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_acosh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_acosh(double x) { @@ -306,12 +306,12 @@ extern "C" __device__ double test_acosh(double x) { // DEFAULT-LABEL: @test_asinf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_asinf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_asin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_asin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_asinf(float x) { @@ -320,12 +320,12 @@ extern "C" __device__ float test_asinf(float x) { // DEFAULT-LABEL: @test_asin( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_asin( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_asin_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_asin_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_asin(double x) { @@ -335,12 +335,12 @@ extern "C" __device__ double test_asin(double x) { // DEFAULT-LABEL: @test_asinhf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_asinhf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_asinh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_asinh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_asinhf(float x) { @@ -349,12 +349,12 @@ extern "C" __device__ float test_asinhf(float x) { // DEFAULT-LABEL: @test_asinh( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_asinh( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_asinh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_asinh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_asinh(double x) { @@ -363,12 +363,12 @@ extern "C" __device__ double test_asinh(double x) { // DEFAULT-LABEL: @test_atan2f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_atan2f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_atan2_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_atan2_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_atan2f(float x, float y) { @@ -377,12 +377,12 @@ extern "C" __device__ float test_atan2f(float x, float y) { // DEFAULT-LABEL: @test_atan2( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_atan2( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_atan2_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_atan2_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_atan2(double x, double y) { @@ -391,12 +391,12 @@ extern "C" __device__ double test_atan2(double x, double y) { // DEFAULT-LABEL: @test_atanf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_atanf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_atan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_atan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_atanf(float x) { @@ -405,12 +405,12 @@ extern "C" __device__ float test_atanf(float x) { // DEFAULT-LABEL: @test_atan( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_atan( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_atan_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_atan_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_atan(double x) { @@ -419,12 +419,12 @@ extern "C" __device__ double test_atan(double x) { // DEFAULT-LABEL: @test_atanhf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_atanhf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_atanh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_atanh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_atanhf(float x) { @@ -433,12 +433,12 @@ extern "C" __device__ float test_atanhf(float x) { // DEFAULT-LABEL: @test_atanh( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_atanh( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_atanh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_atanh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_atanh(double x) { @@ -447,12 +447,12 @@ extern "C" __device__ double test_atanh(double x) { // DEFAULT-LABEL: @test_cbrtf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_cbrtf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cbrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cbrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_cbrtf(float x) { @@ -461,12 +461,12 @@ extern "C" __device__ float test_cbrtf(float x) { // DEFAULT-LABEL: @test_cbrt( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_cbrt( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cbrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cbrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_cbrt(double x) { @@ -475,12 +475,12 @@ extern "C" __device__ double test_cbrt(double x) { // DEFAULT-LABEL: @test_ceilf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_ceilf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ceil_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ceil_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_ceilf(float x) { @@ -489,12 +489,12 @@ extern "C" __device__ float test_ceilf(float x) { // DEFAULT-LABEL: @test_ceil( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_ceil( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ceil_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ceil_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_ceil(double x) { @@ -503,12 +503,12 @@ extern "C" __device__ double test_ceil(double x) { // DEFAULT-LABEL: @test_copysignf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_copysignf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_copysign_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_copysign_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_copysignf(float x, float y) { @@ -517,12 +517,12 @@ extern "C" __device__ float test_copysignf(float x, float y) { // DEFAULT-LABEL: @test_copysign( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_copysign( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_copysign_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_copysign_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_copysign(double x, double y) { @@ -531,12 +531,12 @@ extern "C" __device__ double test_copysign(double x, double y) { // DEFAULT-LABEL: @test_cosf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_cosf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15:[0-9]+]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16:[0-9]+]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_cosf(float x) { @@ -545,12 +545,12 @@ extern "C" __device__ float test_cosf(float x) { // DEFAULT-LABEL: @test_cos( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_cos( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cos_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cos_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_cos(double x) { @@ -559,12 +559,12 @@ extern "C" __device__ double test_cos(double x) { // DEFAULT-LABEL: @test_coshf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_coshf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cosh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cosh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_coshf(float x) { @@ -573,12 +573,12 @@ extern "C" __device__ float test_coshf(float x) { // DEFAULT-LABEL: @test_cosh( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_cosh( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cosh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cosh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_cosh(double x) { @@ -587,12 +587,12 @@ extern "C" __device__ double test_cosh(double x) { // DEFAULT-LABEL: @test_cospif( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_cospif( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cospi_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cospi_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_cospif(float x) { @@ -601,12 +601,12 @@ extern "C" __device__ float test_cospif(float x) { // DEFAULT-LABEL: @test_cospi( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_cospi( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cospi_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cospi_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_cospi(double x) { @@ -615,12 +615,12 @@ extern "C" __device__ double test_cospi(double x) { // DEFAULT-LABEL: @test_cyl_bessel_i0f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_cyl_bessel_i0f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_i0_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_i0_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_cyl_bessel_i0f(float x) { @@ -629,12 +629,12 @@ extern "C" __device__ float test_cyl_bessel_i0f(float x) { // DEFAULT-LABEL: @test_cyl_bessel_i0( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_cyl_bessel_i0( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_i0_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_i0_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_cyl_bessel_i0(double x) { @@ -643,12 +643,12 @@ extern "C" __device__ double test_cyl_bessel_i0(double x) { // DEFAULT-LABEL: @test_cyl_bessel_i1f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_cyl_bessel_i1f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_i1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_i1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_cyl_bessel_i1f(float x) { @@ -657,12 +657,12 @@ extern "C" __device__ float test_cyl_bessel_i1f(float x) { // DEFAULT-LABEL: @test_cyl_bessel_i1( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_cyl_bessel_i1( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_i1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_i1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_cyl_bessel_i1(double x) { @@ -671,12 +671,12 @@ extern "C" __device__ double test_cyl_bessel_i1(double x) { // DEFAULT-LABEL: @test_erfcf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_erfcf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_erfc_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_erfc_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_erfcf(float x) { @@ -685,12 +685,12 @@ extern "C" __device__ float test_erfcf(float x) { // DEFAULT-LABEL: @test_erfc( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_erfc( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_erfc_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_erfc_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_erfc(double x) { @@ -699,12 +699,12 @@ extern "C" __device__ double test_erfc(double x) { // DEFAULT-LABEL: @test_erfinvf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_erfinvf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_erfinv_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_erfinv_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_erfinvf(float x) { @@ -713,12 +713,12 @@ extern "C" __device__ float test_erfinvf(float x) { // DEFAULT-LABEL: @test_erfinv( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_erfinv( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_erfinv_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_erfinv_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_erfinv(double x) { @@ -727,12 +727,12 @@ extern "C" __device__ double test_erfinv(double x) { // DEFAULT-LABEL: @test_exp10f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_exp10f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_exp10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_exp10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_exp10f(float x) { @@ -741,12 +741,12 @@ extern "C" __device__ float test_exp10f(float x) { // DEFAULT-LABEL: @test_exp10( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_exp10( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_exp10_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_exp10_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_exp10(double x) { @@ -755,12 +755,12 @@ extern "C" __device__ double test_exp10(double x) { // DEFAULT-LABEL: @test_exp2f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_exp2f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_exp2_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_exp2_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_exp2f(float x) { @@ -769,12 +769,12 @@ extern "C" __device__ float test_exp2f(float x) { // DEFAULT-LABEL: @test_exp2( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_exp2( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_exp2_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_exp2_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_exp2(double x) { @@ -783,12 +783,12 @@ extern "C" __device__ double test_exp2(double x) { // DEFAULT-LABEL: @test_expf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_expf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_exp_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_exp_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_expf(float x) { @@ -797,12 +797,12 @@ extern "C" __device__ float test_expf(float x) { // DEFAULT-LABEL: @test_exp( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_exp( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_exp_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_exp_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_exp(double x) { @@ -811,12 +811,12 @@ extern "C" __device__ double test_exp(double x) { // DEFAULT-LABEL: @test_expm1f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_expm1f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_expm1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_expm1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_expm1f(float x) { @@ -825,12 +825,12 @@ extern "C" __device__ float test_expm1f(float x) { // DEFAULT-LABEL: @test_expm1( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_expm1( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_expm1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_expm1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_expm1(double x) { @@ -867,12 +867,12 @@ extern "C" __device__ double test_fabs(double x) { // DEFAULT-LABEL: @test_fdimf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_fdimf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fdim_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fdim_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_fdimf(float x, float y) { @@ -881,12 +881,12 @@ extern "C" __device__ float test_fdimf(float x, float y) { // DEFAULT-LABEL: @test_fdim( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_fdim( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fdim_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fdim_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_fdim(double x, double y) { @@ -909,12 +909,12 @@ extern "C" __device__ float test_fdividef(float x, float y) { // DEFAULT-LABEL: @test_floorf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_floorf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_floor_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_floor_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_floorf(float x) { @@ -923,12 +923,12 @@ extern "C" __device__ float test_floorf(float x) { // DEFAULT-LABEL: @test_floor( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_floor( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_floor_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_floor_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_floor(double x) { @@ -979,12 +979,12 @@ extern "C" __device__ double test_fma_rn(double x, double y, double z) { // DEFAULT-LABEL: @test_fmaxf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_fmaxf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmax_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmax_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_fmaxf(float x, float y) { @@ -993,12 +993,12 @@ extern "C" __device__ float test_fmaxf(float x, float y) { // DEFAULT-LABEL: @test_fmax( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_fmax( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmax_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmax_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_fmax(double x, double y) { @@ -1007,12 +1007,12 @@ extern "C" __device__ double test_fmax(double x, double y) { // DEFAULT-LABEL: @test_fminf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_fminf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmin_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmin_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_fminf(float x, float y) { @@ -1021,12 +1021,12 @@ extern "C" __device__ float test_fminf(float x, float y) { // DEFAULT-LABEL: @test_fmin( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_fmin( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmin_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmin_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_fmin(double x, double y) { @@ -1035,12 +1035,12 @@ extern "C" __device__ double test_fmin(double x, double y) { // DEFAULT-LABEL: @test_fmodf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_fmodf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmod_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmod_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_fmodf(float x, float y) { @@ -1049,12 +1049,12 @@ extern "C" __device__ float test_fmodf(float x, float y) { // DEFAULT-LABEL: @test_fmod( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_fmod( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmod_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmod_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_fmod(double x, double y) { @@ -1063,23 +1063,17 @@ extern "C" __device__ double test_fmod(double x, double y) { // DEFAULT-LABEL: @test_frexpf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16:[0-9]+]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_frexp_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] -// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11:![0-9]+]] -// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// DEFAULT-NEXT: ret float [[CALL_I]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f32(float [[X:%.*]]) +// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11:![0-9]+]] +// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract float @llvm.amdgcn.frexp.mant.f32(float [[X]]) +// DEFAULT-NEXT: ret float [[TMP1]] // // FINITEONLY-LABEL: @test_frexpf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) -// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16:[0-9]+]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) float @__ocml_frexp_f32(float noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] -// FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11:![0-9]+]] -// FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]] -// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// FINITEONLY-NEXT: ret float [[CALL_I]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f32(float [[X:%.*]]) +// FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11:![0-9]+]] +// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.frexp.mant.f32(float [[X]]) +// FINITEONLY-NEXT: ret float [[TMP1]] // extern "C" __device__ float test_frexpf(float x, int* y) { return frexpf(x, y); @@ -1087,23 +1081,17 @@ extern "C" __device__ float test_frexpf(float x, int* y) { // DEFAULT-LABEL: @test_frexp( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_frexp_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] -// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f64(double [[X:%.*]]) // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// DEFAULT-NEXT: ret double [[CALL_I]] +// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract double @llvm.amdgcn.frexp.mant.f64(double [[X]]) +// DEFAULT-NEXT: ret double [[TMP1]] // // FINITEONLY-LABEL: @test_frexp( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) -// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) double @__ocml_frexp_f64(double noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] -// FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f64(double [[X:%.*]]) // FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]] -// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// FINITEONLY-NEXT: ret double [[CALL_I]] +// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract double @llvm.amdgcn.frexp.mant.f64(double [[X]]) +// FINITEONLY-NEXT: ret double [[TMP1]] // extern "C" __device__ double test_frexp(double x, int* y) { return frexp(x, y); @@ -1111,12 +1099,12 @@ extern "C" __device__ double test_frexp(double x, int* y) { // DEFAULT-LABEL: @test_hypotf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_hypotf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_hypot_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_hypot_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_hypotf(float x, float y) { @@ -1125,12 +1113,12 @@ extern "C" __device__ float test_hypotf(float x, float y) { // DEFAULT-LABEL: @test_hypot( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_hypot( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_hypot_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_hypot_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_hypot(double x, double y) { @@ -1139,12 +1127,12 @@ extern "C" __device__ double test_hypot(double x, double y) { // DEFAULT-LABEL: @test_ilogbf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret i32 [[CALL_I]] // // FINITEONLY-LABEL: @test_ilogbf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret i32 [[CALL_I]] // extern "C" __device__ int test_ilogbf(float x) { @@ -1153,12 +1141,12 @@ extern "C" __device__ int test_ilogbf(float x) { // DEFAULT-LABEL: @test_ilogb( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret i32 [[CALL_I]] // // FINITEONLY-LABEL: @test_ilogb( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret i32 [[CALL_I]] // extern "C" __device__ int test_ilogb(double x) { @@ -1167,14 +1155,14 @@ extern "C" __device__ int test_ilogb(double x) { // DEFAULT-LABEL: @test___finitef( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // DEFAULT-NEXT: ret i32 [[CONV]] // // FINITEONLY-LABEL: @test___finitef( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // FINITEONLY-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // FINITEONLY-NEXT: ret i32 [[CONV]] @@ -1185,14 +1173,14 @@ extern "C" __device__ BOOL_TYPE test___finitef(float x) { // DEFAULT-LABEL: @test___finite( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // DEFAULT-NEXT: ret i32 [[CONV]] // // FINITEONLY-LABEL: @test___finite( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // FINITEONLY-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // FINITEONLY-NEXT: ret i32 [[CONV]] @@ -1203,14 +1191,14 @@ extern "C" __device__ BOOL_TYPE test___finite(double x) { // DEFAULT-LABEL: @test___isinff( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // DEFAULT-NEXT: ret i32 [[CONV]] // // FINITEONLY-LABEL: @test___isinff( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // FINITEONLY-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // FINITEONLY-NEXT: ret i32 [[CONV]] @@ -1221,14 +1209,14 @@ extern "C" __device__ BOOL_TYPE test___isinff(float x) { // DEFAULT-LABEL: @test___isinf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // DEFAULT-NEXT: ret i32 [[CONV]] // // FINITEONLY-LABEL: @test___isinf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // FINITEONLY-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // FINITEONLY-NEXT: ret i32 [[CONV]] @@ -1239,14 +1227,14 @@ extern "C" __device__ BOOL_TYPE test___isinf(double x) { // DEFAULT-LABEL: @test___isnanf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // DEFAULT-NEXT: ret i32 [[CONV]] // // FINITEONLY-LABEL: @test___isnanf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // FINITEONLY-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // FINITEONLY-NEXT: ret i32 [[CONV]] @@ -1257,14 +1245,14 @@ extern "C" __device__ BOOL_TYPE test___isnanf(float x) { // DEFAULT-LABEL: @test___isnan( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // DEFAULT-NEXT: ret i32 [[CONV]] // // FINITEONLY-LABEL: @test___isnan( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // FINITEONLY-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // FINITEONLY-NEXT: ret i32 [[CONV]] @@ -1275,12 +1263,12 @@ extern "C" __device__ BOOL_TYPE test___isnan(double x) { // DEFAULT-LABEL: @test_j0f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_j0f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_j0f(float x) { @@ -1289,12 +1277,12 @@ extern "C" __device__ float test_j0f(float x) { // DEFAULT-LABEL: @test_j0( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_j0( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_j0(double x) { @@ -1303,12 +1291,12 @@ extern "C" __device__ double test_j0(double x) { // DEFAULT-LABEL: @test_j1f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_j1f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_j1f(float x) { @@ -1317,12 +1305,12 @@ extern "C" __device__ float test_j1f(float x) { // DEFAULT-LABEL: @test_j1( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_j1( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_j1(double x) { @@ -1336,14 +1324,14 @@ extern "C" __device__ double test_j1(double x) { // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // DEFAULT-NEXT: ] // DEFAULT: if.then.i: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] // 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]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] // 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]]) #[[ATTR15]] -// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]] +// 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: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 // DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]] // DEFAULT: for.body.i: @@ -1369,14 +1357,14 @@ extern "C" __device__ double test_j1(double x) { // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // FINITEONLY-NEXT: ] // FINITEONLY: if.then.i: -// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]] // FINITEONLY: if.then2.i: -// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT]] // FINITEONLY: if.end4.i: -// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] -// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // 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: @@ -1406,14 +1394,14 @@ extern "C" __device__ float test_jnf(int x, float y) { // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // DEFAULT-NEXT: ] // DEFAULT: if.then.i: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] // 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]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] // 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]]) #[[ATTR15]] -// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]] +// 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: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 // DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]] // DEFAULT: for.body.i: @@ -1439,14 +1427,14 @@ extern "C" __device__ float test_jnf(int x, float y) { // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // FINITEONLY-NEXT: ] // FINITEONLY: if.then.i: -// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT:%.*]] // FINITEONLY: if.then2.i: -// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT]] // FINITEONLY: if.end4.i: -// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] -// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // 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: @@ -1471,12 +1459,12 @@ extern "C" __device__ double test_jn(int x, double y) { // DEFAULT-LABEL: @test_ldexpf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_ldexpf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ldexp_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ldexp_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_ldexpf(float x, int y) { @@ -1485,12 +1473,12 @@ extern "C" __device__ float test_ldexpf(float x, int y) { // DEFAULT-LABEL: @test_ldexp( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_ldexp( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ldexp_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ldexp_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_ldexp(double x, int y) { @@ -1499,12 +1487,12 @@ extern "C" __device__ double test_ldexp(double x, int y) { // DEFAULT-LABEL: @test_lgammaf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_lgammaf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_lgamma_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_lgamma_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_lgammaf(float x) { @@ -1513,12 +1501,12 @@ extern "C" __device__ float test_lgammaf(float x) { // DEFAULT-LABEL: @test_lgamma( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_lgamma( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_lgamma_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_lgamma_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_lgamma(double x) { @@ -1527,13 +1515,13 @@ extern "C" __device__ double test_lgamma(double x) { // DEFAULT-LABEL: @test_llrintf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -1543,13 +1531,13 @@ extern "C" __device__ long long int test_llrintf(float x) { // DEFAULT-LABEL: @test_llrint( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -1559,13 +1547,13 @@ extern "C" __device__ long long int test_llrint(double x) { // DEFAULT-LABEL: @test_llroundf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -1575,13 +1563,13 @@ extern "C" __device__ long long int test_llroundf(float x) { // DEFAULT-LABEL: @test_llround( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -1591,12 +1579,12 @@ extern "C" __device__ long long int test_llround(double x) { // DEFAULT-LABEL: @test_log10f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_log10f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_log10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_log10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_log10f(float x) { @@ -1605,12 +1593,12 @@ extern "C" __device__ float test_log10f(float x) { // DEFAULT-LABEL: @test_log10( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_log10( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_log10_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_log10_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_log10(double x) { @@ -1619,12 +1607,12 @@ extern "C" __device__ double test_log10(double x) { // DEFAULT-LABEL: @test_log1pf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_log1pf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_log1p_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_log1p_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_log1pf(float x) { @@ -1633,12 +1621,12 @@ extern "C" __device__ float test_log1pf(float x) { // DEFAULT-LABEL: @test_log1p( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_log1p( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_log1p_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_log1p_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_log1p(double x) { @@ -1647,12 +1635,12 @@ extern "C" __device__ double test_log1p(double x) { // DEFAULT-LABEL: @test_log2f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_log2f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_log2_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_log2_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_log2f(float x) { @@ -1661,12 +1649,12 @@ extern "C" __device__ float test_log2f(float x) { // DEFAULT-LABEL: @test_log2( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_log2( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_log2_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_log2_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_log2(double x) { @@ -1675,12 +1663,12 @@ extern "C" __device__ double test_log2(double x) { // DEFAULT-LABEL: @test_logbf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_logbf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_logb_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_logb_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_logbf(float x) { @@ -1689,12 +1677,12 @@ extern "C" __device__ float test_logbf(float x) { // DEFAULT-LABEL: @test_logb( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_logb( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_logb_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_logb_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_logb(double x) { @@ -1703,13 +1691,13 @@ extern "C" __device__ double test_logb(double x) { // DEFAULT-LABEL: @test_lrintf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -1719,13 +1707,13 @@ extern "C" __device__ long int test_lrintf(float x) { // DEFAULT-LABEL: @test_lrint( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -1735,13 +1723,13 @@ extern "C" __device__ long int test_lrint(double x) { // DEFAULT-LABEL: @test_lroundf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -1751,13 +1739,13 @@ extern "C" __device__ long int test_lroundf(float x) { // DEFAULT-LABEL: @test_lround( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -1768,21 +1756,21 @@ extern "C" __device__ long int test_lround(double x) { // DEFAULT-LABEL: @test_modff( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// 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: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15:![0-9]+]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // 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]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) float @__ocml_modf_f32(float noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// 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 nofpclass(nan inf) float @__ocml_modf_f32(float noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15:![0-9]+]] // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]] -// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_modff(float x, float* y) { @@ -1792,21 +1780,21 @@ extern "C" __device__ float test_modff(float x, float* y) { // DEFAULT-LABEL: @test_modf( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// 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: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17:![0-9]+]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // 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]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) double @__ocml_modf_f64(double noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) double @__ocml_modf_f64(double noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17:![0-9]+]] // FINITEONLY-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]] -// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_modf(double x, double* y) { @@ -2044,12 +2032,12 @@ extern "C" __device__ double test_nan_fill() { // DEFAULT-LABEL: @test_nearbyintf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_nearbyintf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_nearbyint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_nearbyint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_nearbyintf(float x) { @@ -2058,12 +2046,12 @@ extern "C" __device__ float test_nearbyintf(float x) { // DEFAULT-LABEL: @test_nearbyint( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_nearbyint( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_nearbyint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_nearbyint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_nearbyint(double x) { @@ -2072,12 +2060,12 @@ extern "C" __device__ double test_nearbyint(double x) { // DEFAULT-LABEL: @test_nextafterf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_nextafterf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_nextafter_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_nextafter_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_nextafterf(float x, float y) { @@ -2086,12 +2074,12 @@ extern "C" __device__ float test_nextafterf(float x, float y) { // DEFAULT-LABEL: @test_nextafter( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_nextafter( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_nextafter_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_nextafter_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_nextafter(double x, double y) { @@ -2100,12 +2088,12 @@ extern "C" __device__ double test_nextafter(double x, double y) { // DEFAULT-LABEL: @test_norm3df( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_norm3df( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_len3_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_len3_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_norm3df(float x, float y, float z) { @@ -2114,12 +2102,12 @@ extern "C" __device__ float test_norm3df(float x, float y, float z) { // DEFAULT-LABEL: @test_norm3d( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_norm3d( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_len3_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_len3_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_norm3d(double x, double y, double z) { @@ -2128,12 +2116,12 @@ extern "C" __device__ double test_norm3d(double x, double y, double z) { // DEFAULT-LABEL: @test_norm4df( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR13]] +// 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: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_norm4df( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_len4_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]], float noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_len4_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]], float noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_norm4df(float x, float y, float z, float w) { @@ -2142,12 +2130,12 @@ extern "C" __device__ float test_norm4df(float x, float y, float z, float w) { // DEFAULT-LABEL: @test_norm4d( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR13]] +// 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: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_norm4d( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_len4_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]], double noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_len4_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]], double noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_norm4d(double x, double y, double z, double w) { @@ -2156,12 +2144,12 @@ extern "C" __device__ double test_norm4d(double x, double y, double z, double w) // DEFAULT-LABEL: @test_normcdff( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_normcdff( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ncdf_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ncdf_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_normcdff(float x) { @@ -2170,12 +2158,12 @@ extern "C" __device__ float test_normcdff(float x) { // DEFAULT-LABEL: @test_normcdf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_normcdf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ncdf_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ncdf_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_normcdf(double x) { @@ -2184,12 +2172,12 @@ extern "C" __device__ double test_normcdf(double x) { // DEFAULT-LABEL: @test_normcdfinvf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_normcdfinvf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ncdfinv_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ncdfinv_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_normcdfinvf(float x) { @@ -2198,12 +2186,12 @@ extern "C" __device__ float test_normcdfinvf(float x) { // DEFAULT-LABEL: @test_normcdfinv( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_normcdfinv( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ncdfinv_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ncdfinv_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_normcdfinv(double x) { @@ -2227,7 +2215,7 @@ extern "C" __device__ double test_normcdfinv(double x) { // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP19:![0-9]+]] // DEFAULT: _ZL5normfiPKf.exit: // DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_normf( @@ -2247,7 +2235,7 @@ extern "C" __device__ double test_normcdfinv(double x) { // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP19:![0-9]+]] // FINITEONLY: _ZL5normfiPKf.exit: // FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sqrt_f32(float noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sqrt_f32(float noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_normf(int x, const float *y) { @@ -2271,7 +2259,7 @@ extern "C" __device__ float test_normf(int x, const float *y) { // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]] // DEFAULT: _ZL4normiPKd.exit: // DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_norm( @@ -2291,7 +2279,7 @@ extern "C" __device__ float test_normf(int x, const float *y) { // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]] // FINITEONLY: _ZL4normiPKd.exit: // FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sqrt_f64(double noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sqrt_f64(double noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_norm(int x, const double *y) { @@ -2300,12 +2288,12 @@ extern "C" __device__ double test_norm(int x, const double *y) { // DEFAULT-LABEL: @test_powf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_powf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_pow_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_pow_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_powf(float x, float y) { @@ -2314,12 +2302,12 @@ extern "C" __device__ float test_powf(float x, float y) { // DEFAULT-LABEL: @test_pow( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_pow( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_pow_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_pow_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_pow(double x, double y) { @@ -2328,12 +2316,12 @@ extern "C" __device__ double test_pow(double x, double y) { // DEFAULT-LABEL: @test_powif( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_powif( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_pown_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_pown_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_powif(float x, int y) { @@ -2342,12 +2330,12 @@ extern "C" __device__ float test_powif(float x, int y) { // DEFAULT-LABEL: @test_powi( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_powi( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_pown_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_pown_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_powi(double x, int y) { @@ -2356,12 +2344,12 @@ extern "C" __device__ double test_powi(double x, int y) { // DEFAULT-LABEL: @test_rcbrtf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_rcbrtf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rcbrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rcbrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_rcbrtf(float x) { @@ -2370,12 +2358,12 @@ extern "C" __device__ float test_rcbrtf(float x) { // DEFAULT-LABEL: @test_rcbrt( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_rcbrt( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rcbrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rcbrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_rcbrt(double x) { @@ -2384,12 +2372,12 @@ extern "C" __device__ double test_rcbrt(double x) { // DEFAULT-LABEL: @test_remainderf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_remainderf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_remainder_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_remainder_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_remainderf(float x, float y) { @@ -2398,12 +2386,12 @@ extern "C" __device__ float test_remainderf(float x, float y) { // DEFAULT-LABEL: @test_remainder( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_remainder( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_remainder_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_remainder_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_remainder(double x, double y) { @@ -2413,21 +2401,21 @@ extern "C" __device__ double test_remainder(double x, double y) { // DEFAULT-LABEL: @test_remquof( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// 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: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]] // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // 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]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) float @__ocml_remquo_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) float @__ocml_remquo_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]] // FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]] -// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_remquof(float x, float y, int* z) { @@ -2437,21 +2425,21 @@ extern "C" __device__ float test_remquof(float x, float y, int* z) { // DEFAULT-LABEL: @test_remquo( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// 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: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]] // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // 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]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) double @__ocml_remquo_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) double @__ocml_remquo_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]] // FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]] -// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_remquo(double x, double y, int* z) { @@ -2460,12 +2448,12 @@ extern "C" __device__ double test_remquo(double x, double y, int* z) { // DEFAULT-LABEL: @test_rhypotf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_rhypotf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rhypot_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rhypot_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_rhypotf(float x, float y) { @@ -2474,12 +2462,12 @@ extern "C" __device__ float test_rhypotf(float x, float y) { // DEFAULT-LABEL: @test_rhypot( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_rhypot( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rhypot_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rhypot_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_rhypot(double x, double y) { @@ -2488,12 +2476,12 @@ extern "C" __device__ double test_rhypot(double x, double y) { // DEFAULT-LABEL: @test_rintf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_rintf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rint_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_rintf(float x) { @@ -2502,12 +2490,12 @@ extern "C" __device__ float test_rintf(float x) { // DEFAULT-LABEL: @test_rint( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_rint( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rint_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_rint(double x) { @@ -2531,7 +2519,7 @@ extern "C" __device__ double test_rint(double x) { // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]] // DEFAULT: _ZL6rnormfiPKf.exit: // DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_rnormf( @@ -2551,7 +2539,7 @@ extern "C" __device__ double test_rint(double x) { // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]] // FINITEONLY: _ZL6rnormfiPKf.exit: // FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rsqrt_f32(float noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rsqrt_f32(float noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_rnormf(int x, const float* y) { @@ -2575,7 +2563,7 @@ extern "C" __device__ float test_rnormf(int x, const float* y) { // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]] // DEFAULT: _ZL5rnormiPKd.exit: // DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_rnorm( @@ -2595,7 +2583,7 @@ extern "C" __device__ float test_rnormf(int x, const float* y) { // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]] // FINITEONLY: _ZL5rnormiPKd.exit: // FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rsqrt_f64(double noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rsqrt_f64(double noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_rnorm(int x, const double* y) { @@ -2604,12 +2592,12 @@ extern "C" __device__ double test_rnorm(int x, const double* y) { // DEFAULT-LABEL: @test_rnorm3df( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_rnorm3df( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rlen3_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rlen3_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_rnorm3df(float x, float y, float z) { @@ -2618,12 +2606,12 @@ extern "C" __device__ float test_rnorm3df(float x, float y, float z) { // DEFAULT-LABEL: @test_rnorm3d( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_rnorm3d( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rlen3_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rlen3_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_rnorm3d(double x, double y, double z) { @@ -2632,12 +2620,12 @@ extern "C" __device__ double test_rnorm3d(double x, double y, double z) { // DEFAULT-LABEL: @test_rnorm4df( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR13]] +// 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: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_rnorm4df( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rlen4_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]], float noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rlen4_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]], float noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_rnorm4df(float x, float y, float z, float w) { @@ -2646,12 +2634,12 @@ extern "C" __device__ float test_rnorm4df(float x, float y, float z, float w) { // DEFAULT-LABEL: @test_rnorm4d( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR13]] +// 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: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_rnorm4d( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rlen4_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]], double noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rlen4_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]], double noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w) { @@ -2660,12 +2648,12 @@ extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w // DEFAULT-LABEL: @test_roundf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_roundf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_roundf(float x) { @@ -2674,12 +2662,12 @@ extern "C" __device__ float test_roundf(float x) { // DEFAULT-LABEL: @test_round( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_round( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_round(double x) { @@ -2688,12 +2676,12 @@ extern "C" __device__ double test_round(double x) { // DEFAULT-LABEL: @test_rsqrtf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_rsqrtf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rsqrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rsqrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_rsqrtf(float x) { @@ -2702,12 +2690,12 @@ extern "C" __device__ float test_rsqrtf(float x) { // DEFAULT-LABEL: @test_rsqrt( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_rsqrt( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rsqrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rsqrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_rsqrt(double x) { @@ -2720,10 +2708,10 @@ extern "C" __device__ double test_rsqrt(double x) { // DEFAULT-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]] // DEFAULT: cond.true.i: // DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32 -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]] // 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) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR14]] // 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]] ] @@ -2735,10 +2723,10 @@ extern "C" __device__ double test_rsqrt(double x) { // FINITEONLY-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]] // FINITEONLY: cond.true.i: // FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32 -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_scalbn_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_scalbn_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]] // FINITEONLY-NEXT: br label [[_ZL8SCALBLNFFL_EXIT:%.*]] // FINITEONLY: cond.false.i: -// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_scalb_f32(float noundef nofpclass(nan inf) [[X]], float noundef nofpclass(nan inf) 0x43E0000000000000) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_scalb_f32(float noundef nofpclass(nan inf) [[X]], float noundef nofpclass(nan inf) 0x43E0000000000000) #[[ATTR14]] // 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]] ] @@ -2754,10 +2742,10 @@ extern "C" __device__ float test_scalblnf(float x, long int y) { // DEFAULT-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]] // DEFAULT: cond.true.i: // DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32 -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]] // 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) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR14]] // 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]] ] @@ -2769,10 +2757,10 @@ extern "C" __device__ float test_scalblnf(float x, long int y) { // FINITEONLY-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]] // FINITEONLY: cond.true.i: // FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32 -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_scalbn_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_scalbn_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]] // FINITEONLY-NEXT: br label [[_ZL7SCALBLNDL_EXIT:%.*]] // FINITEONLY: cond.false.i: -// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_scalb_f64(double noundef nofpclass(nan inf) [[X]], double noundef nofpclass(nan inf) 0x43E0000000000000) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_scalb_f64(double noundef nofpclass(nan inf) [[X]], double noundef nofpclass(nan inf) 0x43E0000000000000) #[[ATTR14]] // 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]] ] @@ -2784,12 +2772,12 @@ extern "C" __device__ double test_scalbln(double x, long int y) { // DEFAULT-LABEL: @test_scalbnf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_scalbnf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_scalbn_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_scalbn_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_scalbnf(float x, int y) { @@ -2798,12 +2786,12 @@ extern "C" __device__ float test_scalbnf(float x, int y) { // DEFAULT-LABEL: @test_scalbn( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_scalbn( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_scalbn_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_scalbn_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_scalbn(double x, int y) { @@ -2812,14 +2800,14 @@ extern "C" __device__ double test_scalbn(double x, int y) { // DEFAULT-LABEL: @test___signbitf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // DEFAULT-NEXT: ret i32 [[CONV]] // // FINITEONLY-LABEL: @test___signbitf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // FINITEONLY-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // FINITEONLY-NEXT: ret i32 [[CONV]] @@ -2830,14 +2818,14 @@ extern "C" __device__ BOOL_TYPE test___signbitf(float x) { // DEFAULT-LABEL: @test___signbit( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // DEFAULT-NEXT: ret i32 [[CONV]] // // FINITEONLY-LABEL: @test___signbit( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // FINITEONLY-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // FINITEONLY-NEXT: ret i32 [[CONV]] @@ -2849,23 +2837,23 @@ extern "C" __device__ BOOL_TYPE test___signbit(double x) { // DEFAULT-LABEL: @test_sincosf( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// 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: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]] // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // 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]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) float @__ocml_sincos_f32(float noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) float @__ocml_sincos_f32(float noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]] // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]] -// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: ret void // extern "C" __device__ void test_sincosf(float x, float *y, float *z) { @@ -2875,23 +2863,23 @@ extern "C" __device__ void test_sincosf(float x, float *y, float *z) { // DEFAULT-LABEL: @test_sincos( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// 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: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]] // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // 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]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) double @__ocml_sincos_f64(double noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) double @__ocml_sincos_f64(double noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]] // FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]] -// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: ret void // extern "C" __device__ void test_sincos(double x, double *y, double *z) { @@ -2901,23 +2889,23 @@ extern "C" __device__ void test_sincos(double x, double *y, double *z) { // DEFAULT-LABEL: @test_sincospif( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// 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: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]] // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]] // DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // 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]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) float @__ocml_sincospi_f32(float noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) float @__ocml_sincospi_f32(float noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]] // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]] -// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: ret void // extern "C" __device__ void test_sincospif(float x, float *y, float *z) { @@ -2927,23 +2915,23 @@ extern "C" __device__ void test_sincospif(float x, float *y, float *z) { // DEFAULT-LABEL: @test_sincospi( // DEFAULT-NEXT: entry: // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) -// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// 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: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]] // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]] // DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]] -// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // 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]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) double @__ocml_sincospi_f64(double noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]] +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract nofpclass(nan inf) double @__ocml_sincospi_f64(double noundef nofpclass(nan inf) [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] // FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]] // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]] // FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]] -// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] // FINITEONLY-NEXT: ret void // extern "C" __device__ void test_sincospi(double x, double *y, double *z) { @@ -2952,12 +2940,12 @@ extern "C" __device__ void test_sincospi(double x, double *y, double *z) { // DEFAULT-LABEL: @test_sinf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_sinf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_sinf(float x) { @@ -2966,12 +2954,12 @@ extern "C" __device__ float test_sinf(float x) { // DEFAULT-LABEL: @test_sin( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_sin( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sin_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sin_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_sin(double x) { @@ -2980,12 +2968,12 @@ extern "C" __device__ double test_sin(double x) { // DEFAULT-LABEL: @test_sinpif( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_sinpif( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sinpi_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sinpi_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_sinpif(float x) { @@ -2994,12 +2982,12 @@ extern "C" __device__ float test_sinpif(float x) { // DEFAULT-LABEL: @test_sinpi( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_sinpi( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sinpi_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sinpi_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_sinpi(double x) { @@ -3008,12 +2996,12 @@ extern "C" __device__ double test_sinpi(double x) { // DEFAULT-LABEL: @test_sqrtf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_sqrtf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sqrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sqrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_sqrtf(float x) { @@ -3022,12 +3010,12 @@ extern "C" __device__ float test_sqrtf(float x) { // DEFAULT-LABEL: @test_sqrt( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_sqrt( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sqrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sqrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_sqrt(double x) { @@ -3036,12 +3024,12 @@ extern "C" __device__ double test_sqrt(double x) { // DEFAULT-LABEL: @test_tanf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_tanf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_tanf(float x) { @@ -3050,12 +3038,12 @@ extern "C" __device__ float test_tanf(float x) { // DEFAULT-LABEL: @test_tan( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_tan( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_tan_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_tan_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_tan(double x) { @@ -3064,12 +3052,12 @@ extern "C" __device__ double test_tan(double x) { // DEFAULT-LABEL: @test_tanhf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_tanhf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tanh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tanh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_tanhf(float x) { @@ -3078,12 +3066,12 @@ extern "C" __device__ float test_tanhf(float x) { // DEFAULT-LABEL: @test_tanh( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_tanh( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_tanh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_tanh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_tanh(double x) { @@ -3092,12 +3080,12 @@ extern "C" __device__ double test_tanh(double x) { // DEFAULT-LABEL: @test_tgammaf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_tgammaf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tgamma_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tgamma_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_tgammaf(float x) { @@ -3106,12 +3094,12 @@ extern "C" __device__ float test_tgammaf(float x) { // DEFAULT-LABEL: @test_tgamma( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_tgamma( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_tgamma_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_tgamma_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_tgamma(double x) { @@ -3120,12 +3108,12 @@ extern "C" __device__ double test_tgamma(double x) { // DEFAULT-LABEL: @test_truncf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_truncf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_trunc_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_trunc_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_truncf(float x) { @@ -3134,12 +3122,12 @@ extern "C" __device__ float test_truncf(float x) { // DEFAULT-LABEL: @test_trunc( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_trunc( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_trunc_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_trunc_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_trunc(double x) { @@ -3148,12 +3136,12 @@ extern "C" __device__ double test_trunc(double x) { // DEFAULT-LABEL: @test_y0f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_y0f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_y0f(float x) { @@ -3162,12 +3150,12 @@ extern "C" __device__ float test_y0f(float x) { // DEFAULT-LABEL: @test_y0( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_y0( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_y0(double x) { @@ -3176,12 +3164,12 @@ extern "C" __device__ double test_y0(double x) { // DEFAULT-LABEL: @test_y1f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_y1f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_y1f(float x) { @@ -3190,12 +3178,12 @@ extern "C" __device__ float test_y1f(float x) { // DEFAULT-LABEL: @test_y1( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_y1( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_y1(double x) { @@ -3209,14 +3197,14 @@ extern "C" __device__ double test_y1(double x) { // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // DEFAULT-NEXT: ] // DEFAULT: if.then.i: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] // 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]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] // 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]]) #[[ATTR15]] -// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]] +// 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: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 // DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]] // DEFAULT: for.body.i: @@ -3242,14 +3230,14 @@ extern "C" __device__ double test_y1(double x) { // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // FINITEONLY-NEXT: ] // FINITEONLY: if.then.i: -// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]] // FINITEONLY: if.then2.i: -// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT]] // FINITEONLY: if.end4.i: -// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] -// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // 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: @@ -3279,14 +3267,14 @@ extern "C" __device__ float test_ynf(int x, float y) { // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // DEFAULT-NEXT: ] // DEFAULT: if.then.i: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] // 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]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] // 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]]) #[[ATTR15]] -// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]] +// 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: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 // DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]] // DEFAULT: for.body.i: @@ -3312,14 +3300,14 @@ extern "C" __device__ float test_ynf(int x, float y) { // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // FINITEONLY-NEXT: ] // FINITEONLY: if.then.i: -// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT:%.*]] // FINITEONLY: if.then2.i: -// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT]] // FINITEONLY: if.end4.i: -// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] -// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // 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: @@ -3344,12 +3332,12 @@ extern "C" __device__ double test_yn(int x, double y) { // DEFAULT-LABEL: @test___cosf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test___cosf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_cos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_cos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___cosf(float x) { @@ -3358,12 +3346,12 @@ extern "C" __device__ float test___cosf(float x) { // DEFAULT-LABEL: @test___exp10f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test___exp10f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_exp10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_exp10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___exp10f(float x) { @@ -3372,12 +3360,12 @@ extern "C" __device__ float test___exp10f(float x) { // DEFAULT-LABEL: @test___expf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test___expf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_exp_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_exp_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___expf(float x) { @@ -3470,12 +3458,12 @@ extern "C" __device__ float test___frsqrt_rn(float x) { // DEFAULT-LABEL: @test___fsqrt_rn( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test___fsqrt_rn( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_sqrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_sqrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fsqrt_rn(float x) { @@ -3498,12 +3486,12 @@ extern "C" __device__ float test___fsub_rn(float x, float y) { // DEFAULT-LABEL: @test___log10f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test___log10f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_log10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_log10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___log10f(float x) { @@ -3512,12 +3500,12 @@ extern "C" __device__ float test___log10f(float x) { // DEFAULT-LABEL: @test___log2f( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test___log2f( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_log2_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_log2_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___log2f(float x) { @@ -3526,12 +3514,12 @@ extern "C" __device__ float test___log2f(float x) { // DEFAULT-LABEL: @test___logf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test___logf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_log_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_log_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___logf(float x) { @@ -3540,12 +3528,12 @@ extern "C" __device__ float test___logf(float x) { // DEFAULT-LABEL: @test___powf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test___powf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_pow_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_pow_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___powf(float x, float y) { @@ -3574,17 +3562,17 @@ extern "C" __device__ float test___saturatef(float x) { // DEFAULT-LABEL: @test___sincosf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]] -// DEFAULT-NEXT: [[CALL1_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL1_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]] // DEFAULT-NEXT: store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]] // DEFAULT-NEXT: ret void // // FINITEONLY-LABEL: @test___sincosf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]] -// FINITEONLY-NEXT: [[CALL1_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_cos_f32(float noundef nofpclass(nan inf) [[X]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL1_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_cos_f32(float noundef nofpclass(nan inf) [[X]]) #[[ATTR16]] // FINITEONLY-NEXT: store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]] // FINITEONLY-NEXT: ret void // @@ -3594,12 +3582,12 @@ extern "C" __device__ void test___sincosf(float x, float *y, float *z) { // DEFAULT-LABEL: @test___sinf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test___sinf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___sinf(float x) { @@ -3608,12 +3596,12 @@ extern "C" __device__ float test___sinf(float x) { // DEFAULT-LABEL: @test___tanf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test___tanf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___tanf(float x) { @@ -3678,12 +3666,12 @@ extern "C" __device__ double test___drcp_rn(double x) { // DEFAULT-LABEL: @test___dsqrt_rn( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test___dsqrt_rn( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sqrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sqrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test___dsqrt_rn(double x) { @@ -3706,12 +3694,12 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) { // DEFAULT-LABEL: @test_float_min( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) float @__ocml_fmin_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmin_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I_I]] // extern "C" __device__ float test_float_min(float x, float y) { @@ -3720,12 +3708,12 @@ extern "C" __device__ float test_float_min(float x, float y) { // DEFAULT-LABEL: @test_float_max( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) float @__ocml_fmax_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmax_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret float [[CALL_I_I]] // extern "C" __device__ float test_float_max(float x, float y) { @@ -3734,12 +3722,12 @@ extern "C" __device__ float test_float_max(float x, float y) { // DEFAULT-LABEL: @test_double_min( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) double @__ocml_fmin_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmin_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I_I]] // extern "C" __device__ double test_double_min(double x, double y) { @@ -3748,12 +3736,12 @@ extern "C" __device__ double test_double_min(double x, double y) { // DEFAULT-LABEL: @test_double_max( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] // 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 nofpclass(nan inf) double @__ocml_fmax_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR13]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmax_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: ret double [[CALL_I_I]] // extern "C" __device__ double test_double_max(double x, double y) { diff --git a/clang/test/Headers/__clang_hip_math_ocml_rounded_ops.hip b/clang/test/Headers/__clang_hip_math_ocml_rounded_ops.hip index 35adb961b8d2d..422b5fb56c4d4 100644 --- a/clang/test/Headers/__clang_hip_math_ocml_rounded_ops.hip +++ b/clang/test/Headers/__clang_hip_math_ocml_rounded_ops.hip @@ -9,7 +9,7 @@ // CHECK-LABEL: @test___fadd_rd( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4:[0-9]+]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3:[0-9]+]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fadd_rd(float x, float y) { @@ -18,7 +18,7 @@ extern "C" __device__ float test___fadd_rd(float x, float y) { // CHECK-LABEL: @test___fadd_rn( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fadd_rn(float x, float y) { @@ -27,7 +27,7 @@ extern "C" __device__ float test___fadd_rn(float x, float y) { // CHECK-LABEL: @test___fadd_ru( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fadd_ru(float x, float y) { @@ -36,7 +36,7 @@ extern "C" __device__ float test___fadd_ru(float x, float y) { // CHECK-LABEL: @test___fadd_rz( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fadd_rz(float x, float y) { @@ -45,7 +45,7 @@ extern "C" __device__ float test___fadd_rz(float x, float y) { // CHECK-LABEL: @test__fmaf_rd( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test__fmaf_rd(float x, float y, float z) { @@ -54,7 +54,7 @@ extern "C" __device__ float test__fmaf_rd(float x, float y, float z) { // CHECK-LABEL: @test__fmaf_rn( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test__fmaf_rn(float x, float y, float z) { @@ -63,7 +63,7 @@ extern "C" __device__ float test__fmaf_rn(float x, float y, float z) { // CHECK-LABEL: @test__fmaf_ru( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test__fmaf_ru(float x, float y, float z) { @@ -72,7 +72,7 @@ extern "C" __device__ float test__fmaf_ru(float x, float y, float z) { // CHECK-LABEL: @test__fmaf_rz( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test__fmaf_rz(float x, float y, float z) { @@ -81,7 +81,7 @@ extern "C" __device__ float test__fmaf_rz(float x, float y, float z) { // CHECK-LABEL: @test___fmul_rd( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fmul_rd(float x, float y) { @@ -90,7 +90,7 @@ extern "C" __device__ float test___fmul_rd(float x, float y) { // CHECK-LABEL: @test___fmul_rn( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fmul_rn(float x, float y) { @@ -99,7 +99,7 @@ extern "C" __device__ float test___fmul_rn(float x, float y) { // CHECK-LABEL: @test___fmul_ru( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fmul_ru(float x, float y) { @@ -108,7 +108,7 @@ extern "C" __device__ float test___fmul_ru(float x, float y) { // CHECK-LABEL: @test___fmul_rz( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fmul_rz(float x, float y) { @@ -117,7 +117,7 @@ extern "C" __device__ float test___fmul_rz(float x, float y) { // CHECK-LABEL: @test___frcp_rd( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtn_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtn_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___frcp_rd(float x) { @@ -126,7 +126,7 @@ extern "C" __device__ float test___frcp_rd(float x) { // CHECK-LABEL: @test___frcp_rn( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rte_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rte_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___frcp_rn(float x) { @@ -135,7 +135,7 @@ extern "C" __device__ float test___frcp_rn(float x) { // CHECK-LABEL: @test___frcp_ru( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtp_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtp_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___frcp_ru(float x) { @@ -144,7 +144,7 @@ extern "C" __device__ float test___frcp_ru(float x) { // CHECK-LABEL: @test___frcp_rz( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtz_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtz_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___frcp_rz(float x) { @@ -153,7 +153,7 @@ extern "C" __device__ float test___frcp_rz(float x) { // CHECK-LABEL: @test___fsqrt_rd( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtn_f32(float noundef [[X:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtn_f32(float noundef [[X:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fsqrt_rd(float x) { @@ -162,7 +162,7 @@ extern "C" __device__ float test___fsqrt_rd(float x) { // CHECK-LABEL: @test___fsqrt_rn( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rte_f32(float noundef [[X:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rte_f32(float noundef [[X:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fsqrt_rn(float x) { @@ -171,7 +171,7 @@ extern "C" __device__ float test___fsqrt_rn(float x) { // CHECK-LABEL: @test___fsqrt_ru( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtp_f32(float noundef [[X:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtp_f32(float noundef [[X:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fsqrt_ru(float x) { @@ -180,7 +180,7 @@ extern "C" __device__ float test___fsqrt_ru(float x) { // CHECK-LABEL: @test___fsqrt_rz( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtz_f32(float noundef [[X:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtz_f32(float noundef [[X:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fsqrt_rz(float x) { @@ -189,7 +189,7 @@ extern "C" __device__ float test___fsqrt_rz(float x) { // CHECK-LABEL: @test___fsub_rd( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fsub_rd(float x, float y) { @@ -198,7 +198,7 @@ extern "C" __device__ float test___fsub_rd(float x, float y) { // CHECK-LABEL: @test___fsub_rn( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fsub_rn(float x, float y) { @@ -207,7 +207,7 @@ extern "C" __device__ float test___fsub_rn(float x, float y) { // CHECK-LABEL: @test___fsub_ru( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fsub_ru(float x, float y) { @@ -216,7 +216,7 @@ extern "C" __device__ float test___fsub_ru(float x, float y) { // CHECK-LABEL: @test___fsub_rz( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test___fsub_rz(float x, float y) { @@ -225,7 +225,7 @@ extern "C" __device__ float test___fsub_rz(float x, float y) { // CHECK-LABEL: @test___dadd_rd( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test___dadd_rd(double x, double y) { @@ -234,7 +234,7 @@ extern "C" __device__ double test___dadd_rd(double x, double y) { // CHECK-LABEL: @test___dadd_rn( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test___dadd_rn(double x, double y) { @@ -243,7 +243,7 @@ extern "C" __device__ double test___dadd_rn(double x, double y) { // CHECK-LABEL: @test___dadd_ru( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test___dadd_ru(double x, double y) { @@ -252,7 +252,7 @@ extern "C" __device__ double test___dadd_ru(double x, double y) { // CHECK-LABEL: @test___dadd_rz( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test___dadd_rz(double x, double y) { @@ -261,7 +261,7 @@ extern "C" __device__ double test___dadd_rz(double x, double y) { // CHECK-LABEL: @test___dmul_rd( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test___dmul_rd(double x, double y) { @@ -270,7 +270,7 @@ extern "C" __device__ double test___dmul_rd(double x, double y) { // CHECK-LABEL: @test___dmul_rn( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test___dmul_rn(double x, double y) { @@ -279,7 +279,7 @@ extern "C" __device__ double test___dmul_rn(double x, double y) { // CHECK-LABEL: @test___dmul_ru( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test___dmul_ru(double x, double y) { @@ -288,7 +288,7 @@ extern "C" __device__ double test___dmul_ru(double x, double y) { // CHECK-LABEL: @test___dmul_rz( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test___dmul_rz(double x, double y) { @@ -298,7 +298,7 @@ extern "C" __device__ double test___dmul_rz(double x, double y) { // CHECK-LABEL: @test___drcp_rd( // CHECK-NEXT: entry: // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtn_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtn_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]] // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float // CHECK-NEXT: ret float [[CONV1]] // @@ -309,7 +309,7 @@ extern "C" __device__ float test___drcp_rd(float x) { // CHECK-LABEL: @test___drcp_rn( // CHECK-NEXT: entry: // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rte_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rte_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]] // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float // CHECK-NEXT: ret float [[CONV1]] // @@ -320,7 +320,7 @@ extern "C" __device__ float test___drcp_rn(float x) { // CHECK-LABEL: @test___drcp_ru( // CHECK-NEXT: entry: // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtp_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtp_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]] // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float // CHECK-NEXT: ret float [[CONV1]] // @@ -331,7 +331,7 @@ extern "C" __device__ float test___drcp_ru(float x) { // CHECK-LABEL: @test___drcp_rz( // CHECK-NEXT: entry: // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtz_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtz_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]] // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float // CHECK-NEXT: ret float [[CONV1]] // @@ -342,7 +342,7 @@ extern "C" __device__ float test___drcp_rz(float x) { // CHECK-LABEL: @test___dsqrt_rd( // CHECK-NEXT: entry: // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtn_f64(double noundef [[CONV]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtn_f64(double noundef [[CONV]]) #[[ATTR3]] // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float // CHECK-NEXT: ret float [[CONV1]] // @@ -353,7 +353,7 @@ extern "C" __device__ float test___dsqrt_rd(float x) { // CHECK-LABEL: @test___dsqrt_rn( // CHECK-NEXT: entry: // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rte_f64(double noundef [[CONV]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rte_f64(double noundef [[CONV]]) #[[ATTR3]] // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float // CHECK-NEXT: ret float [[CONV1]] // @@ -364,7 +364,7 @@ extern "C" __device__ float test___dsqrt_rn(float x) { // CHECK-LABEL: @test___dsqrt_ru( // CHECK-NEXT: entry: // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtp_f64(double noundef [[CONV]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtp_f64(double noundef [[CONV]]) #[[ATTR3]] // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float // CHECK-NEXT: ret float [[CONV1]] // @@ -375,7 +375,7 @@ extern "C" __device__ float test___dsqrt_ru(float x) { // CHECK-LABEL: @test___dsqrt_rz( // CHECK-NEXT: entry: // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtz_f64(double noundef [[CONV]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtz_f64(double noundef [[CONV]]) #[[ATTR3]] // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float // CHECK-NEXT: ret float [[CONV1]] // @@ -385,7 +385,7 @@ extern "C" __device__ float test___dsqrt_rz(float x) { // CHECK-LABEL: @test__fma_rd( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test__fma_rd(double x, double y, double z) { @@ -394,7 +394,7 @@ extern "C" __device__ double test__fma_rd(double x, double y, double z) { // CHECK-LABEL: @test__fma_rn( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test__fma_rn(double x, double y, double z) { @@ -403,7 +403,7 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) { // CHECK-LABEL: @test__fma_ru( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test__fma_ru(double x, double y, double z) { @@ -412,7 +412,7 @@ extern "C" __device__ double test__fma_ru(double x, double y, double z) { // CHECK-LABEL: @test__fma_rz( // CHECK-NEXT: entry: -// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]] +// CHECK-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]] // CHECK-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test__fma_rz(double x, double y, double z) {