Index: clang/test/Headers/__clang_hip_math.hip =================================================================== --- /dev/null +++ clang/test/Headers/__clang_hip_math.hip @@ -0,0 +1,3383 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py +// REQUIRES: amdgpu-registered-target + +// Test without OCML_BASIC_ROUNDED_OPERATIONS +// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \ +// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ +// RUN: -internal-isystem %S/Inputs/include \ +// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ +// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \ +// RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,DEFAULT %s + +// Check that we end up with fast math flags set on intrinsic calls +// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \ +// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ +// RUN: -internal-isystem %S/Inputs/include \ +// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ +// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -ffinite-math-only -o - \ +// RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,FINITEONLY %s + +#define BOOL_TYPE int +typedef unsigned long long uint64_t; + +// CHECK-LABEL: @test___make_mantissa_base8( +// CHECK-NEXT: entry: +// CHECK-NEXT: br label [[WHILE_COND_I:%.*]] +// CHECK: while.cond.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ] +// CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_1_I:%.*]], [[CLEANUP_I]] ] +// CHECK-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq ptr [[__TAGP_ADDR_0_I]], null +// CHECK-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// CHECK: while.body.i: +// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA3:![0-9]+]] +// CHECK-NEXT: [[TMP1:%.*]] = and i8 [[TMP0]], -8 +// CHECK-NEXT: [[TMP2:%.*]] = icmp eq i8 [[TMP1]], 48 +// CHECK-NEXT: br i1 [[TMP2]], label [[IF_THEN_I:%.*]], label [[CLEANUP_I]] +// CHECK: if.then.i: +// CHECK-NEXT: [[MUL_I:%.*]] = shl i64 [[__R_0_I]], 3 +// CHECK-NEXT: [[CONV3_I:%.*]] = sext i8 [[TMP0]] to i64 +// CHECK-NEXT: [[ADD_I:%.*]] = add i64 [[MUL_I]], -48 +// CHECK-NEXT: [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV3_I]] +// CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I]] +// CHECK: cleanup.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ], [ [[__TAGP_ADDR_0_I]], [[WHILE_BODY_I]] ] +// CHECK-NEXT: [[__R_1_I]] = phi i64 [ [[SUB_I]], [[IF_THEN_I]] ], [ [[__R_0_I]], [[WHILE_BODY_I]] ] +// CHECK-NEXT: br i1 [[TMP2]], label [[WHILE_COND_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT]], !llvm.loop [[LOOP6:![0-9]+]] +// CHECK: _ZL21__make_mantissa_base8PKc.exit: +// CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ] +// CHECK-NEXT: ret i64 [[RETVAL_2_I]] +// +extern "C" __device__ uint64_t test___make_mantissa_base8(const char *p) { + return __make_mantissa_base8(p); +} + +// CHECK-LABEL: @test___make_mantissa_base10( +// CHECK-NEXT: entry: +// CHECK-NEXT: br label [[WHILE_COND_I:%.*]] +// CHECK: while.cond.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ] +// CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_1_I:%.*]], [[CLEANUP_I]] ] +// CHECK-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq ptr [[__TAGP_ADDR_0_I]], null +// CHECK-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// CHECK: while.body.i: +// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[TMP1:%.*]] = add i8 [[TMP0]], -48 +// CHECK-NEXT: [[TMP2:%.*]] = icmp ult i8 [[TMP1]], 10 +// CHECK-NEXT: br i1 [[TMP2]], label [[IF_THEN_I:%.*]], label [[CLEANUP_I]] +// CHECK: if.then.i: +// CHECK-NEXT: [[MUL_I:%.*]] = mul i64 [[__R_0_I]], 10 +// CHECK-NEXT: [[CONV3_I:%.*]] = sext i8 [[TMP0]] to i64 +// CHECK-NEXT: [[ADD_I:%.*]] = add i64 [[MUL_I]], -48 +// CHECK-NEXT: [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV3_I]] +// CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I]] +// CHECK: cleanup.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ], [ [[__TAGP_ADDR_0_I]], [[WHILE_BODY_I]] ] +// CHECK-NEXT: [[__R_1_I]] = phi i64 [ [[SUB_I]], [[IF_THEN_I]] ], [ [[__R_0_I]], [[WHILE_BODY_I]] ] +// CHECK-NEXT: br i1 [[TMP2]], label [[WHILE_COND_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT]], !llvm.loop [[LOOP9:![0-9]+]] +// CHECK: _ZL22__make_mantissa_base10PKc.exit: +// CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ] +// CHECK-NEXT: ret i64 [[RETVAL_2_I]] +// +extern "C" __device__ uint64_t test___make_mantissa_base10(const char *p) { + return __make_mantissa_base10(p); +} + +// CHECK-LABEL: @test___make_mantissa_base16( +// CHECK-NEXT: entry: +// CHECK-NEXT: br label [[WHILE_COND_I:%.*]] +// CHECK: while.cond.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ] +// CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_2_I:%.*]], [[CLEANUP_I]] ] +// CHECK-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq ptr [[__TAGP_ADDR_0_I]], null +// CHECK-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// CHECK: while.body.i: +// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[TMP1:%.*]] = add i8 [[TMP0]], -48 +// CHECK-NEXT: [[TMP2:%.*]] = icmp ult i8 [[TMP1]], 10 +// CHECK-NEXT: br i1 [[TMP2]], label [[IF_END29_I:%.*]], label [[IF_ELSE_I:%.*]] +// CHECK: if.else.i: +// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP0]], -97 +// CHECK-NEXT: [[TMP4:%.*]] = icmp ult i8 [[TMP3]], 6 +// CHECK-NEXT: br i1 [[TMP4]], label [[IF_END29_I]], label [[IF_ELSE15_I:%.*]] +// CHECK: if.else15.i: +// CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP0]], -65 +// CHECK-NEXT: [[TMP6:%.*]] = icmp ult i8 [[TMP5]], 6 +// CHECK-NEXT: br i1 [[TMP6]], label [[IF_END29_I]], label [[CLEANUP_I]] +// CHECK: if.end29.i: +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I]] ], [ -87, [[IF_ELSE_I]] ], [ -55, [[IF_ELSE15_I]] ] +// CHECK-NEXT: [[MUL22_I:%.*]] = shl i64 [[__R_0_I]], 4 +// CHECK-NEXT: [[CONV23_I:%.*]] = sext i8 [[TMP0]] to i64 +// CHECK-NEXT: [[ADD24_I:%.*]] = add i64 [[MUL22_I]], [[DOTSINK]] +// CHECK-NEXT: [[ADD26_I:%.*]] = add i64 [[ADD24_I]], [[CONV23_I]] +// CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I]] +// CHECK: cleanup.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_END29_I]] ], [ [[__TAGP_ADDR_0_I]], [[IF_ELSE15_I]] ] +// CHECK-NEXT: [[__R_2_I]] = phi i64 [ [[ADD26_I]], [[IF_END29_I]] ], [ [[__R_0_I]], [[IF_ELSE15_I]] ] +// CHECK-NEXT: [[COND_I:%.*]] = phi i1 [ true, [[IF_END29_I]] ], [ false, [[IF_ELSE15_I]] ] +// CHECK-NEXT: br i1 [[COND_I]], label [[WHILE_COND_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]], !llvm.loop [[LOOP10:![0-9]+]] +// CHECK: _ZL22__make_mantissa_base16PKc.exit: +// CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ] +// CHECK-NEXT: ret i64 [[RETVAL_2_I]] +// +extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) { + return __make_mantissa_base16(p); +} + +// CHECK-LABEL: @test___make_mantissa( +// CHECK-NEXT: entry: +// CHECK-NEXT: ret i64 0 +// +extern "C" __device__ uint64_t test___make_mantissa(const char *p) { + return __make_mantissa(p); +} + +// CHECK-LABEL: @test_abs( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.abs.i32(i32 [[X:%.*]], i1 true) +// CHECK-NEXT: ret i32 [[TMP0]] +// +extern "C" __device__ int test_abs(int x) { + return abs(x); +} + +// CHECK-LABEL: @test_labs( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[TMP0:%.*]] = tail call i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true) +// CHECK-NEXT: ret i64 [[TMP0]] +// +extern "C" __device__ long test_labs(long x) { + return labs(x); +} + +// +// CHECK-LABEL: @test_llabs( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[TMP0:%.*]] = tail call i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true) +// CHECK-NEXT: ret i64 [[TMP0]] +// +extern "C" __device__ long long test_llabs(long x) { + return llabs(x); +} + +// DEFAULT-LABEL: @test_acosf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_acosf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_acosf(float x) { + return acosf(x); +} + +// DEFAULT-LABEL: @test_acos( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_acos( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_acos(double x) { + return acos(x); +} + +// DEFAULT-LABEL: @test_acoshf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_acoshf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_acoshf(float x) { + return acoshf(x); +} + +// DEFAULT-LABEL: @test_acosh( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_acosh( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_acosh(double x) { + return acosh(x); +} + +// DEFAULT-LABEL: @test_asinf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_asinf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_asinf(float x) { + return asinf(x); +} + +// DEFAULT-LABEL: @test_asin( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_asin( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_asin(double x) { + + return asin(x); +} + +// DEFAULT-LABEL: @test_asinhf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_asinhf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_asinhf(float x) { + return asinhf(x); +} + +// DEFAULT-LABEL: @test_asinh( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_asinh( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_asinh(double x) { + return asinh(x); +} + +// DEFAULT-LABEL: @test_atan2f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_atan2f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_atan2f(float x, float y) { + return atan2f(x, y); +} + +// DEFAULT-LABEL: @test_atan2( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_atan2( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_atan2(double x, double y) { + return atan2(x, y); +} + +// DEFAULT-LABEL: @test_atanf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_atanf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_atanf(float x) { + return atanf(x); +} + +// DEFAULT-LABEL: @test_atan( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_atan( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_atan(double x) { + return atan(x); +} + +// DEFAULT-LABEL: @test_atanhf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_atanhf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_atanhf(float x) { + return atanhf(x); +} + +// DEFAULT-LABEL: @test_atanh( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_atanh( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_atanh(double x) { + return atanh(x); +} + +// DEFAULT-LABEL: @test_cbrtf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_cbrtf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_cbrtf(float x) { + return cbrtf(x); +} + +// DEFAULT-LABEL: @test_cbrt( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_cbrt( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_cbrt(double x) { + return cbrt(x); +} + +// DEFAULT-LABEL: @test_ceilf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_ceilf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_ceilf(float x) { + return ceilf(x); +} + +// DEFAULT-LABEL: @test_ceil( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_ceil( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_ceil(double x) { + return ceil(x); +} + +// DEFAULT-LABEL: @test_copysignf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_copysignf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_copysignf(float x, float y) { + return copysignf(x, y); +} + +// DEFAULT-LABEL: @test_copysign( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_copysign( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_copysign(double x, double y) { + return copysign(x, y); +} + +// DEFAULT-LABEL: @test_cosf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_cosf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_cosf(float x) { + return cosf(x); +} + +// DEFAULT-LABEL: @test_cos( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_cos( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_cos(double x) { + return cos(x); +} + +// DEFAULT-LABEL: @test_coshf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_coshf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_coshf(float x) { + return coshf(x); +} + +// DEFAULT-LABEL: @test_cosh( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_cosh( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_cosh(double x) { + return cosh(x); +} + +// DEFAULT-LABEL: @test_cospif( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_cospif( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_cospif(float x) { + return cospif(x); +} + +// DEFAULT-LABEL: @test_cospi( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_cospi( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_cospi(double x) { + return cospi(x); +} + +// DEFAULT-LABEL: @test_cyl_bessel_i0f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_cyl_bessel_i0f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_cyl_bessel_i0f(float x) { + return cyl_bessel_i0f(x); +} + +// DEFAULT-LABEL: @test_cyl_bessel_i0( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_cyl_bessel_i0( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_cyl_bessel_i0(double x) { + return cyl_bessel_i0(x); +} + +// DEFAULT-LABEL: @test_cyl_bessel_i1f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_cyl_bessel_i1f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_cyl_bessel_i1f(float x) { + return cyl_bessel_i1f(x); +} + +// DEFAULT-LABEL: @test_cyl_bessel_i1( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_cyl_bessel_i1( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_cyl_bessel_i1(double x) { + return cyl_bessel_i1(x); +} + +// DEFAULT-LABEL: @test_erfcf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_erfcf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_erfcf(float x) { + return erfcf(x); +} + +// DEFAULT-LABEL: @test_erfc( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_erfc( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_erfc(double x) { + return erfc(x); +} + +// DEFAULT-LABEL: @test_erfinvf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_erfinvf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_erfinvf(float x) { + return erfinvf(x); +} + +// DEFAULT-LABEL: @test_erfinv( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_erfinv( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_erfinv(double x) { + return erfinv(x); +} + +// DEFAULT-LABEL: @test_exp10f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_exp10f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_exp10f(float x) { + return exp10f(x); +} + +// DEFAULT-LABEL: @test_exp10( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_exp10( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_exp10(double x) { + return exp10(x); +} + +// DEFAULT-LABEL: @test_exp2f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_exp2f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_exp2f(float x) { + return exp2f(x); +} + +// DEFAULT-LABEL: @test_exp2( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_exp2( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_exp2(double x) { + return exp2(x); +} + +// DEFAULT-LABEL: @test_expf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_expf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_expf(float x) { + return expf(x); +} + +// DEFAULT-LABEL: @test_exp( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_exp( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_exp(double x) { + return exp(x); +} + +// DEFAULT-LABEL: @test_expm1f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_expm1f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_expm1f(float x) { + return expm1f(x); +} + +// DEFAULT-LABEL: @test_expm1( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_expm1( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_expm1(double x) { + return expm1(x); +} + +// DEFAULT-LABEL: @test_fabsf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fabs_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fabsf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fabs_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_fabsf(float x) { + return fabsf(x); +} + +// DEFAULT-LABEL: @test_fabs( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fabs_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fabs( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fabs_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_fabs(double x) { + return fabs(x); +} + +// DEFAULT-LABEL: @test_fdimf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fdimf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_fdimf(float x, float y) { + return fdimf(x, y); +} + +// DEFAULT-LABEL: @test_fdim( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fdim( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_fdim(double x, double y) { + return fdim(x, y); +} + +// DEFAULT-LABEL: @test_fdividef( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]] +// DEFAULT-NEXT: ret float [[DIV_I]] +// +// FINITEONLY-LABEL: @test_fdividef( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[X:%.*]], [[Y:%.*]] +// FINITEONLY-NEXT: ret float [[DIV_I]] +// +extern "C" __device__ float test_fdividef(float x, float y) { + return fdividef(x, y); +} + +// DEFAULT-LABEL: @test_floorf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_floorf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_floorf(float x) { + return floorf(x); +} + +// DEFAULT-LABEL: @test_floor( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_floor( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_floor(double x) { + return floor(x); +} + +// DEFAULT-LABEL: @test_fmaf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fmaf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_fmaf(float x, float y, float z) { + return fmaf(x, y, z); +} + +// DEFAULT-LABEL: @test_fma( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fma( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_fma(double x, double y, double z) { + return fma(x, y, z); +} + +// DEFAULT-LABEL: @test_fma_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fma_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_fma_rn(double x, double y, double z) { + return __fma_rn(x, y, z); +} + +// DEFAULT-LABEL: @test_fmaxf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fmaxf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_fmaxf(float x, float y) { + return fmaxf(x, y); +} + +// DEFAULT-LABEL: @test_fmax( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fmax( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_fmax(double x, double y) { + return fmax(x, y); +} + +// DEFAULT-LABEL: @test_fminf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fminf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_fminf(float x, float y) { + return fminf(x, y); +} + +// DEFAULT-LABEL: @test_fmin( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fmin( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_fmin(double x, double y) { + return fmin(x, y); +} + +// DEFAULT-LABEL: @test_fmodf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fmodf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_fmodf(float x, float y) { + return fmodf(x, y); +} + +// DEFAULT-LABEL: @test_fmod( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_fmod( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_fmod(double x, double y) { + return fmod(x, y); +} + +// DEFAULT-LABEL: @test_hypotf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_hypotf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_hypotf(float x, float y) { + return hypotf(x, y); +} + +// DEFAULT-LABEL: @test_hypot( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_hypot( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_hypot(double x, double y) { + return hypot(x, y); +} + +// CHECK-LABEL: @test_ilogbf( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]] +// CHECK-NEXT: ret i32 [[CALL_I]] +// +extern "C" __device__ int test_ilogbf(float x) { + return ilogbf(x); +} + +// CHECK-LABEL: @test_ilogb( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// CHECK-NEXT: ret i32 [[CALL_I]] +// +extern "C" __device__ int test_ilogb(double x) { + return ilogb(x); +} + +// CHECK-LABEL: @test___finitef( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 +// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 +// CHECK-NEXT: ret i32 [[CONV]] +// +extern "C" __device__ BOOL_TYPE test___finitef(float x) { + return __finitef(x); +} + +// CHECK-LABEL: @test___finite( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 +// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 +// CHECK-NEXT: ret i32 [[CONV]] +// +extern "C" __device__ BOOL_TYPE test___finite(double x) { + return __finite(x); +} + +// CHECK-LABEL: @test___isinff( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 +// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 +// CHECK-NEXT: ret i32 [[CONV]] +// +extern "C" __device__ BOOL_TYPE test___isinff(float x) { + return __isinff(x); +} + +// CHECK-LABEL: @test___isinf( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 +// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 +// CHECK-NEXT: ret i32 [[CONV]] +// +extern "C" __device__ BOOL_TYPE test___isinf(double x) { + return __isinf(x); +} + +// CHECK-LABEL: @test___isnanf( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 +// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 +// CHECK-NEXT: ret i32 [[CONV]] +// +extern "C" __device__ BOOL_TYPE test___isnanf(float x) { + return __isnanf(x); +} + +// CHECK-LABEL: @test___isnan( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 +// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 +// CHECK-NEXT: ret i32 [[CONV]] +// +extern "C" __device__ BOOL_TYPE test___isnan(double x) { + return __isnan(x); +} + +// DEFAULT-LABEL: @test_j0f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_j0f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_j0f(float x) { + return j0f(x); +} + +// DEFAULT-LABEL: @test_j0( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_j0( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_j0(double x) { + return j0(x); +} + +// DEFAULT-LABEL: @test_j1f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_j1f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_j1f(float x) { + return j1f(x); +} + +// DEFAULT-LABEL: @test_j1( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_j1( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_j1(double x) { + return j1(x); +} + +// DEFAULT-LABEL: @test_jnf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [ +// DEFAULT-NEXT: i32 0, label [[IF_THEN_I:%.*]] +// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] +// DEFAULT-NEXT: ] +// DEFAULT: if.then.i: +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]] +// DEFAULT: if.then2.i: +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT]] +// DEFAULT: if.end4.i: +// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 +// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]] +// DEFAULT: for.body.i: +// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float +// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]] +// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]] +// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]] +// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP11:![0-9]+]] +// DEFAULT: _ZL3jnfif.exit: +// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// DEFAULT-NEXT: ret float [[RETVAL_0_I]] +// +// FINITEONLY-LABEL: @test_jnf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [ +// FINITEONLY-NEXT: i32 0, label [[IF_THEN_I:%.*]] +// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] +// FINITEONLY-NEXT: ] +// FINITEONLY: if.then.i: +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]] +// FINITEONLY: if.then2.i: +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT]] +// FINITEONLY: if.end4.i: +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 +// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]] +// FINITEONLY: for.body.i: +// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float +// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]] +// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]] +// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]] +// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP11:![0-9]+]] +// FINITEONLY: _ZL3jnfif.exit: +// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// FINITEONLY-NEXT: ret float [[RETVAL_0_I]] +// +extern "C" __device__ float test_jnf(int x, float y) { + return jnf(x, y); +} + +// DEFAULT-LABEL: @test_jn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [ +// DEFAULT-NEXT: i32 0, label [[IF_THEN_I:%.*]] +// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] +// DEFAULT-NEXT: ] +// DEFAULT: if.then.i: +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: br label [[_ZL2JNID_EXIT:%.*]] +// DEFAULT: if.then2.i: +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: br label [[_ZL2JNID_EXIT]] +// DEFAULT: if.end4.i: +// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 +// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]] +// DEFAULT: for.body.i: +// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double +// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]] +// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]] +// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]] +// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP12:![0-9]+]] +// DEFAULT: _ZL2jnid.exit: +// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// DEFAULT-NEXT: ret double [[RETVAL_0_I]] +// +// FINITEONLY-LABEL: @test_jn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [ +// FINITEONLY-NEXT: i32 0, label [[IF_THEN_I:%.*]] +// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] +// FINITEONLY-NEXT: ] +// FINITEONLY: if.then.i: +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT:%.*]] +// FINITEONLY: if.then2.i: +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT]] +// FINITEONLY: if.end4.i: +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 +// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]] +// FINITEONLY: for.body.i: +// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double +// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]] +// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]] +// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]] +// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP12:![0-9]+]] +// FINITEONLY: _ZL2jnid.exit: +// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// FINITEONLY-NEXT: ret double [[RETVAL_0_I]] +// +extern "C" __device__ double test_jn(int x, double y) { + return jn(x, y); +} + +// DEFAULT-LABEL: @test_ldexpf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_ldexpf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_ldexpf(float x, int y) { + return ldexpf(x, y); +} + +// DEFAULT-LABEL: @test_ldexp( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_ldexp( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_ldexp(double x, int y) { + return ldexp(x, y); +} + +// DEFAULT-LABEL: @test_lgammaf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_lgammaf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_lgammaf(float x) { + return lgammaf(x); +} + +// DEFAULT-LABEL: @test_lgamma( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_lgamma( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_lgamma(double x) { + return lgamma(x); +} + +// DEFAULT-LABEL: @test_llrintf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// DEFAULT-NEXT: ret i64 [[CONV_I]] +// +// FINITEONLY-LABEL: @test_llrintf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// FINITEONLY-NEXT: ret i64 [[CONV_I]] +// +extern "C" __device__ long long int test_llrintf(float x) { + return llrintf(x); +} + +// DEFAULT-LABEL: @test_llrint( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// DEFAULT-NEXT: ret i64 [[CONV_I]] +// +// FINITEONLY-LABEL: @test_llrint( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// FINITEONLY-NEXT: ret i64 [[CONV_I]] +// +extern "C" __device__ long long int test_llrint(double x) { + return llrint(x); +} + +// DEFAULT-LABEL: @test_llroundf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// DEFAULT-NEXT: ret i64 [[CONV_I]] +// +// FINITEONLY-LABEL: @test_llroundf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// FINITEONLY-NEXT: ret i64 [[CONV_I]] +// +extern "C" __device__ long long int test_llroundf(float x) { + return llroundf(x); +} + +// DEFAULT-LABEL: @test_llround( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// DEFAULT-NEXT: ret i64 [[CONV_I]] +// +// FINITEONLY-LABEL: @test_llround( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// FINITEONLY-NEXT: ret i64 [[CONV_I]] +// +extern "C" __device__ long long int test_llround(double x) { + return llround(x); +} + +// DEFAULT-LABEL: @test_log10f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_log10f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_log10f(float x) { + return log10f(x); +} + +// DEFAULT-LABEL: @test_log10( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_log10( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_log10(double x) { + return log10(x); +} + +// DEFAULT-LABEL: @test_log1pf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_log1pf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_log1pf(float x) { + return log1pf(x); +} + +// DEFAULT-LABEL: @test_log1p( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_log1p( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_log1p(double x) { + return log1p(x); +} + +// DEFAULT-LABEL: @test_log2f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_log2f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_log2f(float x) { + return log2f(x); +} + +// DEFAULT-LABEL: @test_log2( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_log2( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_log2(double x) { + return log2(x); +} + +// DEFAULT-LABEL: @test_logbf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_logbf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_logbf(float x) { + return logbf(x); +} + +// DEFAULT-LABEL: @test_logb( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_logb( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_logb(double x) { + return logb(x); +} + +// DEFAULT-LABEL: @test_lrintf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// DEFAULT-NEXT: ret i64 [[CONV_I]] +// +// FINITEONLY-LABEL: @test_lrintf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// FINITEONLY-NEXT: ret i64 [[CONV_I]] +// +extern "C" __device__ long int test_lrintf(float x) { + return lrintf(x); +} + +// DEFAULT-LABEL: @test_lrint( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// DEFAULT-NEXT: ret i64 [[CONV_I]] +// +// FINITEONLY-LABEL: @test_lrint( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// FINITEONLY-NEXT: ret i64 [[CONV_I]] +// +extern "C" __device__ long int test_lrint(double x) { + return lrint(x); +} + +// DEFAULT-LABEL: @test_lroundf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// DEFAULT-NEXT: ret i64 [[CONV_I]] +// +// FINITEONLY-LABEL: @test_lroundf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 +// FINITEONLY-NEXT: ret i64 [[CONV_I]] +// +extern "C" __device__ long int test_lroundf(float x) { + return lroundf(x); +} + +// DEFAULT-LABEL: @test_lround( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// DEFAULT-NEXT: ret i64 [[CONV_I]] +// +// FINITEONLY-LABEL: @test_lround( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 +// FINITEONLY-NEXT: ret i64 [[CONV_I]] +// +extern "C" __device__ long int test_lround(double x) { + return lround(x); +} + +// DEFAULT-LABEL: @test_modff( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17:[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 [[TBAA13:![0-9]+]] +// DEFAULT-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_modff( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17:[0-9]+]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13:![0-9]+]] +// FINITEONLY-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_modff(float x, float* y) { + return modff(x, y); +} + +// DEFAULT-LABEL: @test_modf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15:![0-9]+]] +// DEFAULT-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_modf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15:![0-9]+]] +// FINITEONLY-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_modf(double x, double* y) { + return modf(x, y); +} + +// CHECK-LABEL: @test_nanf( +// CHECK-NEXT: entry: +// CHECK-NEXT: ret float 0x7FF8000000000000 +// +extern "C" __device__ float test_nanf(const char *tag) { + return nanf(tag); +} + +// CHECK-LABEL: @test_nan( +// CHECK-NEXT: entry: +// CHECK-NEXT: ret double 0x7FF8000000000000 +// +extern "C" __device__ double test_nan(const char *tag) { + return nan(tag); +} + +// CHECK-LABEL: @test_nanf_emptystr( +// CHECK-NEXT: entry: +// CHECK-NEXT: ret float 0x7FF8000000000000 +// +extern "C" __device__ float test_nanf_emptystr() { + return nanf(""); +} + +// CHECK-LABEL: @test_nan_emptystr( +// CHECK-NEXT: entry: +// CHECK-NEXT: ret double 0x7FF8000000000000 +// +extern "C" __device__ double test_nan_emptystr() { + return nan(""); +} + +// CHECK-LABEL: @test_nanf_fill( +// CHECK-NEXT: entry: +// CHECK-NEXT: ret float 0x7FF8000000000000 +// +extern "C" __device__ float test_nanf_fill() { + return nanf("0x456"); +} + +// CHECK-LABEL: @test_nan_fill( +// CHECK-NEXT: entry: +// CHECK-NEXT: ret double 0x7FF8000000000000 +// +extern "C" __device__ double test_nan_fill() { + return nan("0x123"); +} + +// DEFAULT-LABEL: @test_nearbyintf( +// DEFAULT-NEXT: entry: +// 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 float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_nearbyintf(float x) { + return nearbyintf(x); +} + +// DEFAULT-LABEL: @test_nearbyint( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_nearbyint( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_nearbyint(double x) { + return nearbyint(x); +} + +// DEFAULT-LABEL: @test_nextafterf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_nextafterf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_nextafterf(float x, float y) { + return nextafterf(x, y); +} + +// DEFAULT-LABEL: @test_nextafter( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_nextafter( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_nextafter(double x, double y) { + return nextafter(x, y); +} + +// DEFAULT-LABEL: @test_norm3df( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_norm3df( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_norm3df(float x, float y, float z) { + return norm3df(x, y, z); +} + +// DEFAULT-LABEL: @test_norm3d( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_norm3d( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_norm3d(double x, double y, double z) { + return norm3d(x, y, z); +} + +// DEFAULT-LABEL: @test_norm4df( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_norm4df( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_norm4df(float x, float y, float z, float w) { + return norm4df(x, y, z, w); +} + +// DEFAULT-LABEL: @test_norm4d( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_norm4d( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_norm4d(double x, double y, double z, double w) { + return norm4d(x, y, z, w); +} + +// DEFAULT-LABEL: @test_normcdff( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_normcdff( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_normcdff(float x) { + return normcdff(x); +} + +// DEFAULT-LABEL: @test_normcdf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_normcdf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_normcdf(double x) { + return normcdf(x); +} + +// DEFAULT-LABEL: @test_normcdfinvf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_normcdfinvf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_normcdfinvf(float x) { + return normcdfinvf(x); +} + +// DEFAULT-LABEL: @test_normcdfinv( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_normcdfinv( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_normcdfinv(double x) { + return normcdfinv(x); +} + +// DEFAULT-LABEL: @test_normf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// DEFAULT: while.body.i: +// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 +// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA13]] +// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]] +// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]] +// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1 +// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP17:![0-9]+]] +// DEFAULT: _ZL5normfiPKf.exit: +// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_normf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// FINITEONLY: while.body.i: +// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 +// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA13]] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]] +// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]] +// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1 +// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP17:![0-9]+]] +// FINITEONLY: _ZL5normfiPKf.exit: +// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_normf(int x, const float *y) { + return normf(x, y); +} + +// DEFAULT-LABEL: @test_norm( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// DEFAULT: while.body.i: +// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 +// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA15]] +// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]] +// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]] +// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1 +// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP18:![0-9]+]] +// DEFAULT: _ZL4normiPKd.exit: +// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_norm( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// FINITEONLY: while.body.i: +// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 +// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA15]] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]] +// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]] +// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1 +// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP18:![0-9]+]] +// FINITEONLY: _ZL4normiPKd.exit: +// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_norm(int x, const double *y) { + return norm(x, y); +} + +// DEFAULT-LABEL: @test_powf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_powf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_powf(float x, float y) { + return powf(x, y); +} + +// DEFAULT-LABEL: @test_pow( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_pow( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_pow(double x, double y) { + return pow(x, y); +} + +// DEFAULT-LABEL: @test_powif( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_powif( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_powif(float x, int y) { + return powif(x, y); +} + +// DEFAULT-LABEL: @test_powi( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_powi( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_powi(double x, int y) { + return powi(x, y); +} + +// DEFAULT-LABEL: @test_rcbrtf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rcbrtf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_rcbrtf(float x) { + return rcbrtf(x); +} + +// DEFAULT-LABEL: @test_rcbrt( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rcbrt( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_rcbrt(double x) { + return rcbrt(x); +} + +// DEFAULT-LABEL: @test_remainderf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_remainderf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_remainderf(float x, float y) { + return remainderf(x, y); +} + +// DEFAULT-LABEL: @test_remainder( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_remainder( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_remainder(double x, double y) { + return remainder(x, y); +} + +// DEFAULT-LABEL: @test_remquof( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19:![0-9]+]] +// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_remquof( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19:![0-9]+]] +// FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_remquof(float x, float y, int* z) { + return remquof(x, y, z); +} + +// DEFAULT-LABEL: @test_remquo( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19]] +// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_remquo( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5) +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19]] +// FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_remquo(double x, double y, int* z) { + return remquo(x, y, z); +} + +// DEFAULT-LABEL: @test_rhypotf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rhypotf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_rhypotf(float x, float y) { + return rhypotf(x, y); +} + +// DEFAULT-LABEL: @test_rhypot( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rhypot( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_rhypot(double x, double y) { + return rhypot(x, y); +} + +// DEFAULT-LABEL: @test_rintf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rintf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_rintf(float x) { + return rintf(x); +} + +// DEFAULT-LABEL: @test_rint( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rint( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_rint(double x) { + return rint(x); +} + +// DEFAULT-LABEL: @test_rnormf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// DEFAULT: while.body.i: +// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 +// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA13]] +// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]] +// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]] +// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1 +// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]] +// DEFAULT: _ZL6rnormfiPKf.exit: +// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rnormf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// FINITEONLY: while.body.i: +// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 +// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA13]] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]] +// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]] +// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1 +// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]] +// FINITEONLY: _ZL6rnormfiPKf.exit: +// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_rnormf(int x, const float* y) { + return rnormf(x, y); +} + +// DEFAULT-LABEL: @test_rnorm( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// DEFAULT: while.body.i: +// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 +// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA15]] +// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]] +// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]] +// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1 +// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]] +// DEFAULT: _ZL5rnormiPKd.exit: +// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rnorm( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// FINITEONLY: while.body.i: +// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 +// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA15]] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]] +// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]] +// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1 +// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]] +// FINITEONLY: _ZL5rnormiPKd.exit: +// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_rnorm(int x, const double* y) { + return rnorm(x, y); +} + +// DEFAULT-LABEL: @test_rnorm3df( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rnorm3df( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_rnorm3df(float x, float y, float z) { + return rnorm3df(x, y, z); +} + +// DEFAULT-LABEL: @test_rnorm3d( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rnorm3d( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_rnorm3d(double x, double y, double z) { + return rnorm3d(x, y, z); +} + +// DEFAULT-LABEL: @test_rnorm4df( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rnorm4df( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_rnorm4df(float x, float y, float z, float w) { + return rnorm4df(x, y, z, w); +} + +// DEFAULT-LABEL: @test_rnorm4d( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rnorm4d( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w) { + return rnorm4d(x, y, z, w); +} + +// DEFAULT-LABEL: @test_roundf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_roundf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_roundf(float x) { + return roundf(x); +} + +// DEFAULT-LABEL: @test_round( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_round( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_round(double x) { + return round(x); +} + +// DEFAULT-LABEL: @test_rsqrtf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rsqrtf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_rsqrtf(float x) { + return rsqrtf(x); +} + +// DEFAULT-LABEL: @test_rsqrt( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_rsqrt( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_rsqrt(double x) { + return rsqrt(x); +} + +// DEFAULT-LABEL: @test_scalblnf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807 +// DEFAULT-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]] +// DEFAULT: cond.true.i: +// DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32 +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]] +// DEFAULT-NEXT: br label [[_ZL8SCALBLNFFL_EXIT:%.*]] +// DEFAULT: cond.false.i: +// DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR14]] +// DEFAULT-NEXT: br label [[_ZL8SCALBLNFFL_EXIT]] +// DEFAULT: _ZL8scalblnffl.exit: +// DEFAULT-NEXT: [[COND_I:%.*]] = phi contract float [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ] +// DEFAULT-NEXT: ret float [[COND_I]] +// +// FINITEONLY-LABEL: @test_scalblnf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807 +// FINITEONLY-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]] +// FINITEONLY: cond.true.i: +// FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32 +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]] +// FINITEONLY-NEXT: br label [[_ZL8SCALBLNFFL_EXIT:%.*]] +// FINITEONLY: cond.false.i: +// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR14]] +// FINITEONLY-NEXT: 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]] ] +// FINITEONLY-NEXT: ret float [[COND_I]] +// +extern "C" __device__ float test_scalblnf(float x, long int y) { + return scalblnf(x, y); +} + +// DEFAULT-LABEL: @test_scalbln( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807 +// DEFAULT-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]] +// DEFAULT: cond.true.i: +// DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32 +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]] +// DEFAULT-NEXT: br label [[_ZL7SCALBLNDL_EXIT:%.*]] +// DEFAULT: cond.false.i: +// DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR14]] +// DEFAULT-NEXT: br label [[_ZL7SCALBLNDL_EXIT]] +// DEFAULT: _ZL7scalblndl.exit: +// DEFAULT-NEXT: [[COND_I:%.*]] = phi contract double [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ] +// DEFAULT-NEXT: ret double [[COND_I]] +// +// FINITEONLY-LABEL: @test_scalbln( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807 +// FINITEONLY-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]] +// FINITEONLY: cond.true.i: +// FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32 +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]] +// FINITEONLY-NEXT: br label [[_ZL7SCALBLNDL_EXIT:%.*]] +// FINITEONLY: cond.false.i: +// FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR14]] +// FINITEONLY-NEXT: 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]] ] +// FINITEONLY-NEXT: ret double [[COND_I]] +// +extern "C" __device__ double test_scalbln(double x, long int y) { + return scalbln(x, y); +} + +// DEFAULT-LABEL: @test_scalbnf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_scalbnf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_scalbnf(float x, int y) { + return scalbnf(x, y); +} + +// DEFAULT-LABEL: @test_scalbn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_scalbn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_scalbn(double x, int y) { + return scalbn(x, y); +} + +// CHECK-LABEL: @test___signbitf( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 +// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 +// CHECK-NEXT: ret i32 [[CONV]] +// +extern "C" __device__ BOOL_TYPE test___signbitf(float x) { + return __signbitf(x); +} + +// CHECK-LABEL: @test___signbit( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 +// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 +// CHECK-NEXT: ret i32 [[CONV]] +// +extern "C" __device__ BOOL_TYPE test___signbit(double x) { + return __signbit(x); +} + +// DEFAULT-LABEL: @test_sincosf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] +// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] +// DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: ret void +// +// FINITEONLY-LABEL: @test_sincosf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] +// FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: ret void +// +extern "C" __device__ void test_sincosf(float x, float *y, float *z) { + sincosf(x, y, z); +} + +// DEFAULT-LABEL: @test_sincos( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] +// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]] +// DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: ret void +// +// FINITEONLY-LABEL: @test_sincos( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]] +// FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: ret void +// +extern "C" __device__ void test_sincos(double x, double *y, double *z) { + sincos(x, y, z); +} + +// DEFAULT-LABEL: @test_sincospif( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] +// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] +// DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: ret void +// +// FINITEONLY-LABEL: @test_sincospif( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]] +// FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: ret void +// +extern "C" __device__ void test_sincospif(float x, float *y, float *z) { + sincospif(x, y, z); +} + +// DEFAULT-LABEL: @test_sincospi( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) +// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] +// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]] +// DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]] +// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// DEFAULT-NEXT: ret void +// +// FINITEONLY-LABEL: @test_sincospi( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5) +// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]] +// FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]] +// FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]] +// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]] +// FINITEONLY-NEXT: ret void +// +extern "C" __device__ void test_sincospi(double x, double *y, double *z) { + sincospi(x, y, z); +} + +// DEFAULT-LABEL: @test_sinf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_sinf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_sinf(float x) { + return sinf(x); +} + +// DEFAULT-LABEL: @test_sin( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_sin( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_sin(double x) { + return sin(x); +} + +// DEFAULT-LABEL: @test_sinpif( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_sinpif( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_sinpif(float x) { + return sinpif(x); +} + +// DEFAULT-LABEL: @test_sinpi( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_sinpi( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_sinpi(double x) { + return sinpi(x); +} + +// DEFAULT-LABEL: @test_sqrtf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_sqrtf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_sqrtf(float x) { + return sqrtf(x); +} + +// DEFAULT-LABEL: @test_sqrt( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_sqrt( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_sqrt(double x) { + return sqrt(x); +} + +// DEFAULT-LABEL: @test_tanf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_tanf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_tanf(float x) { + return tanf(x); +} + +// DEFAULT-LABEL: @test_tan( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_tan( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_tan(double x) { + return tan(x); +} + +// DEFAULT-LABEL: @test_tanhf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_tanhf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_tanhf(float x) { + return tanhf(x); +} + +// DEFAULT-LABEL: @test_tanh( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_tanh( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_tanh(double x) { + return tanh(x); +} + +// DEFAULT-LABEL: @test_tgammaf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_tgammaf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_tgammaf(float x) { + return tgammaf(x); +} + +// DEFAULT-LABEL: @test_tgamma( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_tgamma( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_tgamma(double x) { + return tgamma(x); +} + +// DEFAULT-LABEL: @test_truncf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_truncf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_truncf(float x) { + return truncf(x); +} + +// DEFAULT-LABEL: @test_trunc( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_trunc( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_trunc(double x) { + return trunc(x); +} + +// DEFAULT-LABEL: @test_y0f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_y0f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_y0f(float x) { + return y0f(x); +} + +// DEFAULT-LABEL: @test_y0( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_y0( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_y0(double x) { + return y0(x); +} + +// DEFAULT-LABEL: @test_y1f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test_y1f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test_y1f(float x) { + return y1f(x); +} + +// DEFAULT-LABEL: @test_y1( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test_y1( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test_y1(double x) { + return y1(x); +} + +// DEFAULT-LABEL: @test_ynf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [ +// DEFAULT-NEXT: i32 0, label [[IF_THEN_I:%.*]] +// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] +// DEFAULT-NEXT: ] +// DEFAULT: if.then.i: +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]] +// DEFAULT: if.then2.i: +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT]] +// DEFAULT: if.end4.i: +// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 +// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]] +// DEFAULT: for.body.i: +// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float +// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]] +// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]] +// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]] +// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]] +// DEFAULT: _ZL3ynfif.exit: +// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// DEFAULT-NEXT: ret float [[RETVAL_0_I]] +// +// FINITEONLY-LABEL: @test_ynf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [ +// FINITEONLY-NEXT: i32 0, label [[IF_THEN_I:%.*]] +// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] +// FINITEONLY-NEXT: ] +// FINITEONLY: if.then.i: +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]] +// FINITEONLY: if.then2.i: +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT]] +// FINITEONLY: if.end4.i: +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 +// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]] +// FINITEONLY: for.body.i: +// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float +// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]] +// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]] +// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]] +// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]] +// FINITEONLY: _ZL3ynfif.exit: +// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// FINITEONLY-NEXT: ret float [[RETVAL_0_I]] +// +extern "C" __device__ float test_ynf(int x, float y) { + return ynf(x, y); +} + +// DEFAULT-LABEL: @test_yn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [ +// DEFAULT-NEXT: i32 0, label [[IF_THEN_I:%.*]] +// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] +// DEFAULT-NEXT: ] +// DEFAULT: if.then.i: +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: br label [[_ZL2YNID_EXIT:%.*]] +// DEFAULT: if.then2.i: +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: br label [[_ZL2YNID_EXIT]] +// DEFAULT: if.end4.i: +// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 +// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]] +// DEFAULT: for.body.i: +// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double +// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]] +// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]] +// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]] +// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]] +// DEFAULT: _ZL2ynid.exit: +// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// DEFAULT-NEXT: ret double [[RETVAL_0_I]] +// +// FINITEONLY-LABEL: @test_yn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [ +// FINITEONLY-NEXT: i32 0, label [[IF_THEN_I:%.*]] +// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] +// FINITEONLY-NEXT: ] +// FINITEONLY: if.then.i: +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT:%.*]] +// FINITEONLY: if.then2.i: +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT]] +// FINITEONLY: if.end4.i: +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 +// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]] +// FINITEONLY: for.body.i: +// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double +// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]] +// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]] +// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]] +// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]] +// FINITEONLY: _ZL2ynid.exit: +// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// FINITEONLY-NEXT: ret double [[RETVAL_0_I]] +// +extern "C" __device__ double test_yn(int x, double y) { + return yn(x, y); +} + +// DEFAULT-LABEL: @test___cosf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test___cosf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test___cosf(float x) { + return __cosf(x); +} + +// DEFAULT-LABEL: @test___exp10f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test___exp10f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test___exp10f(float x) { + return __exp10f(x); +} + +// DEFAULT-LABEL: @test___expf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test___expf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test___expf(float x) { + return __expf(x); +} + +// DEFAULT-LABEL: @test___fadd_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[ADD_I:%.*]] = fadd contract float [[X:%.*]], [[Y:%.*]] +// DEFAULT-NEXT: ret float [[ADD_I]] +// +// FINITEONLY-LABEL: @test___fadd_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[ADD_I:%.*]] = fadd nnan ninf contract float [[X:%.*]], [[Y:%.*]] +// FINITEONLY-NEXT: ret float [[ADD_I]] +// +extern "C" __device__ float test___fadd_rn(float x, float y) { + return __fadd_rn(x, y); +} + +// DEFAULT-LABEL: @test___fdividef( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]] +// DEFAULT-NEXT: ret float [[DIV_I]] +// +// FINITEONLY-LABEL: @test___fdividef( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[X:%.*]], [[Y:%.*]] +// FINITEONLY-NEXT: ret float [[DIV_I]] +// +extern "C" __device__ float test___fdividef(float x, float y) { + return __fdividef(x, y); +} + +// DEFAULT-LABEL: @test__fmaf_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test__fmaf_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test__fmaf_rn(float x, float y, float z) { + return __fmaf_rn(x, y, z); +} + +// DEFAULT-LABEL: @test___fmul_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[X:%.*]], [[Y:%.*]] +// DEFAULT-NEXT: ret float [[MUL_I]] +// +// FINITEONLY-LABEL: @test___fmul_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[X:%.*]], [[Y:%.*]] +// FINITEONLY-NEXT: ret float [[MUL_I]] +// +extern "C" __device__ float test___fmul_rn(float x, float y) { + return __fmul_rn(x, y); +} + +// DEFAULT-LABEL: @test___frcp_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float 1.000000e+00, [[X:%.*]] +// DEFAULT-NEXT: ret float [[DIV_I]] +// +// FINITEONLY-LABEL: @test___frcp_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float 1.000000e+00, [[X:%.*]] +// FINITEONLY-NEXT: ret float [[DIV_I]] +// +extern "C" __device__ float test___frcp_rn(float x) { + return __frcp_rn(x); +} + +// DEFAULT-LABEL: @test___fsqrt_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test___fsqrt_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test___fsqrt_rn(float x) { + return __fsqrt_rn(x); +} + +// DEFAULT-LABEL: @test___fsub_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[SUB_I:%.*]] = fsub contract float [[X:%.*]], [[Y:%.*]] +// DEFAULT-NEXT: ret float [[SUB_I]] +// +// FINITEONLY-LABEL: @test___fsub_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[SUB_I:%.*]] = fsub nnan ninf contract float [[X:%.*]], [[Y:%.*]] +// FINITEONLY-NEXT: ret float [[SUB_I]] +// +extern "C" __device__ float test___fsub_rn(float x, float y) { + return __fsub_rn(x, y); +} + +// DEFAULT-LABEL: @test___log10f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test___log10f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test___log10f(float x) { + return __log10f(x); +} + +// DEFAULT-LABEL: @test___log2f( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test___log2f( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test___log2f(float x) { + return __log2f(x); +} + +// DEFAULT-LABEL: @test___logf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test___logf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test___logf(float x) { + return __logf(x); +} + +// DEFAULT-LABEL: @test___powf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test___powf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test___powf(float x, float y) { + return __powf(x, y); +} + +// DEFAULT-LABEL: @test___saturatef( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CMP_I:%.*]] = fcmp contract olt float [[X:%.*]], 0.000000e+00 +// DEFAULT-NEXT: [[CMP1_I:%.*]] = fcmp contract ogt float [[X]], 1.000000e+00 +// DEFAULT-NEXT: [[COND_I:%.*]] = select contract i1 [[CMP1_I]], float 1.000000e+00, float [[X]] +// DEFAULT-NEXT: [[COND5_I:%.*]] = select contract i1 [[CMP_I]], float 0.000000e+00, float [[COND_I]] +// DEFAULT-NEXT: ret float [[COND5_I]] +// +// FINITEONLY-LABEL: @test___saturatef( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CMP_I:%.*]] = fcmp nnan ninf contract olt float [[X:%.*]], 0.000000e+00 +// FINITEONLY-NEXT: [[CMP1_I:%.*]] = fcmp nnan ninf contract ogt float [[X]], 1.000000e+00 +// FINITEONLY-NEXT: [[COND_I:%.*]] = select nnan ninf contract i1 [[CMP1_I]], float 1.000000e+00, float [[X]] +// FINITEONLY-NEXT: [[COND5_I:%.*]] = select nnan ninf contract i1 [[CMP_I]], float 0.000000e+00, float [[COND_I]] +// FINITEONLY-NEXT: ret float [[COND5_I]] +// +extern "C" __device__ float test___saturatef(float x) { + return __saturatef(x); +} + +// DEFAULT-LABEL: @test___sincosf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] +// DEFAULT-NEXT: [[CALL1_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]] +// DEFAULT-NEXT: store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] +// DEFAULT-NEXT: ret void +// +// FINITEONLY-LABEL: @test___sincosf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]] +// FINITEONLY-NEXT: [[CALL1_I:%.*]] = tail call nnan ninf contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]] +// FINITEONLY-NEXT: store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]] +// FINITEONLY-NEXT: ret void +// +extern "C" __device__ void test___sincosf(float x, float *y, float *z) { + __sincosf(x, y, z); +} + +// DEFAULT-LABEL: @test___sinf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test___sinf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test___sinf(float x) { + return __sinf(x); +} + +// DEFAULT-LABEL: @test___tanf( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: ret float [[CALL_I]] +// +// FINITEONLY-LABEL: @test___tanf( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: ret float [[CALL_I]] +// +extern "C" __device__ float test___tanf(float x) { + return __tanf(x); +} + +// DEFAULT-LABEL: @test___dadd_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[ADD_I:%.*]] = fadd contract double [[X:%.*]], [[Y:%.*]] +// DEFAULT-NEXT: ret double [[ADD_I]] +// +// FINITEONLY-LABEL: @test___dadd_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[ADD_I:%.*]] = fadd nnan ninf contract double [[X:%.*]], [[Y:%.*]] +// FINITEONLY-NEXT: ret double [[ADD_I]] +// +extern "C" __device__ double test___dadd_rn(double x, double y) { + return __dadd_rn(x, y); +} + +// DEFAULT-LABEL: @test___ddiv_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[X:%.*]], [[Y:%.*]] +// DEFAULT-NEXT: ret double [[DIV_I]] +// +// FINITEONLY-LABEL: @test___ddiv_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[X:%.*]], [[Y:%.*]] +// FINITEONLY-NEXT: ret double [[DIV_I]] +// +extern "C" __device__ double test___ddiv_rn(double x, double y) { + return __ddiv_rn(x, y); +} + +// DEFAULT-LABEL: @test___dmul_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[X:%.*]], [[Y:%.*]] +// DEFAULT-NEXT: ret double [[MUL_I]] +// +// FINITEONLY-LABEL: @test___dmul_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[X:%.*]], [[Y:%.*]] +// FINITEONLY-NEXT: ret double [[MUL_I]] +// +extern "C" __device__ double test___dmul_rn(double x, double y) { + return __dmul_rn(x, y); +} + +// DEFAULT-LABEL: @test___drcp_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double 1.000000e+00, [[X:%.*]] +// DEFAULT-NEXT: ret double [[DIV_I]] +// +// FINITEONLY-LABEL: @test___drcp_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double 1.000000e+00, [[X:%.*]] +// FINITEONLY-NEXT: ret double [[DIV_I]] +// +extern "C" __device__ double test___drcp_rn(double x) { + return __drcp_rn(x); +} + +// DEFAULT-LABEL: @test___dsqrt_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test___dsqrt_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test___dsqrt_rn(double x) { + return __dsqrt_rn(x); +} + +// DEFAULT-LABEL: @test__fma_rn( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I]] +// +// FINITEONLY-LABEL: @test__fma_rn( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I]] +// +extern "C" __device__ double test__fma_rn(double x, double y, double z) { + return __fma_rn(x, y, z); +} + +// DEFAULT-LABEL: @test_float_min( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I_I]] +// +// FINITEONLY-LABEL: @test_float_min( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I_I]] +// +extern "C" __device__ float test_float_min(float x, float y) { + return min(x, y); +} + +// DEFAULT-LABEL: @test_float_max( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret float [[CALL_I_I]] +// +// FINITEONLY-LABEL: @test_float_max( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret float [[CALL_I_I]] +// +extern "C" __device__ float test_float_max(float x, float y) { + return max(x, y); +} + +// DEFAULT-LABEL: @test_double_min( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I_I]] +// +// FINITEONLY-LABEL: @test_double_min( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I_I]] +// +extern "C" __device__ double test_double_min(double x, double y) { + return min(x, y); +} + +// DEFAULT-LABEL: @test_double_max( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// DEFAULT-NEXT: ret double [[CALL_I_I]] +// +// FINITEONLY-LABEL: @test_double_max( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]] +// FINITEONLY-NEXT: ret double [[CALL_I_I]] +// +extern "C" __device__ double test_double_max(double x, double y) { + return max(x, y); +} +// CHECK-LABEL: @test_int_min( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.smin.i32(i32 [[X:%.*]], i32 [[Y:%.*]]) +// CHECK-NEXT: ret i32 [[TMP0]] +// +extern "C" __device__ int test_int_min(int x, int y) { + return min(x, y); +} + +// CHECK-LABEL: @test_int_max( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.smax.i32(i32 [[X:%.*]], i32 [[Y:%.*]]) +// CHECK-NEXT: ret i32 [[TMP0]] +// +extern "C" __device__ int test_int_max(int x, int y) { + return max(x, y); +} Index: clang/test/Headers/__clang_hip_math_ocml_rounded_ops.hip =================================================================== --- /dev/null +++ clang/test/Headers/__clang_hip_math_ocml_rounded_ops.hip @@ -0,0 +1,420 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py +// Test with OCML_BASIC_ROUNDED_OPERATIONS +// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \ +// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ +// RUN: -internal-isystem %S/Inputs/include \ +// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ +// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \ +// RUN: -D__HIPCC_RTC__ -DOCML_BASIC_ROUNDED_OPERATIONS | FileCheck %s + +// 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fadd_rd(float x, float y) { + return __fadd_rd(x, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fadd_rn(float x, float y) { + return __fadd_rn(x, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fadd_ru(float x, float y) { + return __fadd_ru(x, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fadd_rz(float x, float y) { + return __fadd_rz(x, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test__fmaf_rd(float x, float y, float z) { + return __fmaf_rd(x, y, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test__fmaf_rn(float x, float y, float z) { + return __fmaf_rn(x, y, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test__fmaf_ru(float x, float y, float z) { + return __fmaf_ru(x, y, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test__fmaf_rz(float x, float y, float z) { + return __fmaf_rz(x, y, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fmul_rd(float x, float y) { + return __fmul_rd(x, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fmul_rn(float x, float y) { + return __fmul_rn(x, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fmul_ru(float x, float y) { + return __fmul_ru(x, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fmul_rz(float x, float y) { + return __fmul_rz(x, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___frcp_rd(float x) { + return __frcp_rd(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: ret float [[CALL_I]] +// +extern "C" __device__ float test___frcp_rn(float x) { + return __frcp_rn(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: ret float [[CALL_I]] +// +extern "C" __device__ float test___frcp_ru(float x) { + return __frcp_ru(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: ret float [[CALL_I]] +// +extern "C" __device__ float test___frcp_rz(float x) { + return __frcp_rz(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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fsqrt_rd(float x) { + return __fsqrt_rd(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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fsqrt_rn(float x) { + return __fsqrt_rn(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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fsqrt_ru(float x) { + return __fsqrt_ru(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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fsqrt_rz(float x) { + return __fsqrt_rz(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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fsub_rd(float x, float y) { + return __fsub_rd(x, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fsub_rn(float x, float y) { + return __fsub_rn(x, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fsub_ru(float x, float y) { + return __fsub_ru(x, 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: ret float [[CALL_I]] +// +extern "C" __device__ float test___fsub_rz(float x, float y) { + return __fsub_rz(x, 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: ret double [[CALL_I]] +// +extern "C" __device__ double test___dadd_rd(double x, double y) { + return __dadd_rd(x, 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: ret double [[CALL_I]] +// +extern "C" __device__ double test___dadd_rn(double x, double y) { + return __dadd_rn(x, 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: ret double [[CALL_I]] +// +extern "C" __device__ double test___dadd_ru(double x, double y) { + return __dadd_ru(x, 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: ret double [[CALL_I]] +// +extern "C" __device__ double test___dadd_rz(double x, double y) { + return __dadd_rz(x, 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: ret double [[CALL_I]] +// +extern "C" __device__ double test___dmul_rd(double x, double y) { + return __dmul_rd(x, 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: ret double [[CALL_I]] +// +extern "C" __device__ double test___dmul_rn(double x, double y) { + return __dmul_rn(x, 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: ret double [[CALL_I]] +// +extern "C" __device__ double test___dmul_ru(double x, double y) { + return __dmul_ru(x, 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: ret double [[CALL_I]] +// +extern "C" __device__ double test___dmul_rz(double x, double y) { + return __dmul_rz(x, 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: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float +// CHECK-NEXT: ret float [[CONV1]] +// +extern "C" __device__ float test___drcp_rd(float x) { + return __drcp_rd(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: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float +// CHECK-NEXT: ret float [[CONV1]] +// +extern "C" __device__ float test___drcp_rn(float x) { + return __drcp_rn(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: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float +// CHECK-NEXT: ret float [[CONV1]] +// +extern "C" __device__ float test___drcp_ru(float x) { + return __drcp_ru(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: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float +// CHECK-NEXT: ret float [[CONV1]] +// +extern "C" __device__ float test___drcp_rz(float x) { + return __drcp_rz(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: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float +// CHECK-NEXT: ret float [[CONV1]] +// +extern "C" __device__ float test___dsqrt_rd(float x) { + return __dsqrt_rd(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: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float +// CHECK-NEXT: ret float [[CONV1]] +// +extern "C" __device__ float test___dsqrt_rn(float x) { + return __dsqrt_rn(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: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float +// CHECK-NEXT: ret float [[CONV1]] +// +extern "C" __device__ float test___dsqrt_ru(float x) { + return __dsqrt_ru(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: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float +// CHECK-NEXT: ret float [[CONV1]] +// +extern "C" __device__ float test___dsqrt_rz(float x) { + return __dsqrt_rz(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: ret double [[CALL_I]] +// +extern "C" __device__ double test__fma_rd(double x, double y, double z) { + return __fma_rd(x, y, 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: ret double [[CALL_I]] +// +extern "C" __device__ double test__fma_rn(double x, double y, double z) { + return __fma_rn(x, y, 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: ret double [[CALL_I]] +// +extern "C" __device__ double test__fma_ru(double x, double y, double z) { + return __fma_ru(x, y, 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: ret double [[CALL_I]] +// +extern "C" __device__ double test__fma_rz(double x, double y, double z) { + return __fma_rz(x, y, z); +}