Index: clang/lib/Headers/__clang_hip_math.h =================================================================== --- clang/lib/Headers/__clang_hip_math.h +++ clang/lib/Headers/__clang_hip_math.h @@ -70,9 +70,9 @@ #endif __DEVICE__ -uint64_t __make_mantissa_base8(const char *__tagp) { +uint64_t __make_mantissa_base8(const char *__tagp __attribute__((nonnull))) { uint64_t __r = 0; - while (__tagp) { + while (*__tagp != '\0') { char __tmp = *__tagp; if (__tmp >= '0' && __tmp <= '7') @@ -87,9 +87,9 @@ } __DEVICE__ -uint64_t __make_mantissa_base10(const char *__tagp) { +uint64_t __make_mantissa_base10(const char *__tagp __attribute__((nonnull))) { uint64_t __r = 0; - while (__tagp) { + while (*__tagp != '\0') { char __tmp = *__tagp; if (__tmp >= '0' && __tmp <= '9') @@ -104,9 +104,9 @@ } __DEVICE__ -uint64_t __make_mantissa_base16(const char *__tagp) { +uint64_t __make_mantissa_base16(const char *__tagp __attribute__((nonnull))) { uint64_t __r = 0; - while (__tagp) { + while (*__tagp != '\0') { char __tmp = *__tagp; if (__tmp >= '0' && __tmp <= '9') @@ -125,10 +125,7 @@ } __DEVICE__ -uint64_t __make_mantissa(const char *__tagp) { - if (!__tagp) - return 0u; - +uint64_t __make_mantissa(const char *__tagp __attribute__((nonnull))) { if (*__tagp == '0') { ++__tagp; @@ -359,7 +356,7 @@ } __DEVICE__ -float nanf(const char *__tagp) { +float nanf(const char *__tagp __attribute__((nonnull))) { union { float val; struct ieee_float { Index: clang/test/Headers/__clang_hip_math.hip =================================================================== --- clang/test/Headers/__clang_hip_math.hip +++ clang/test/Headers/__clang_hip_math.hip @@ -26,18 +26,18 @@ // 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: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// CHECK: while.body.i: // 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: [[CONV5_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: [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV5_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: @@ -58,18 +58,18 @@ // 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: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// CHECK: while.body.i: // 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: [[CONV5_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: [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV5_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: @@ -84,39 +84,40 @@ 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: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// CHECK: while.body.i: // 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-NEXT: br i1 [[TMP2]], label [[IF_END31_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: br i1 [[TMP4]], label [[IF_END31_I]], label [[IF_ELSE17_I:%.*]] +// CHECK: if.else17.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: br i1 [[TMP6]], label [[IF_END31_I]], label [[CLEANUP_I]] +// CHECK: if.end31.i: +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I]] ], [ -87, [[IF_ELSE_I]] ], [ -55, [[IF_ELSE17_I]] ] +// CHECK-NEXT: [[MUL24_I:%.*]] = shl i64 [[__R_0_I]], 4 +// CHECK-NEXT: [[CONV25_I:%.*]] = sext i8 [[TMP0]] to i64 +// CHECK-NEXT: [[ADD26_I:%.*]] = add i64 [[MUL24_I]], [[DOTSINK]] +// CHECK-NEXT: [[ADD28_I:%.*]] = add i64 [[ADD26_I]], [[CONV25_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: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_END31_I]] ], [ [[__TAGP_ADDR_0_I]], [[IF_ELSE17_I]] ] +// CHECK-NEXT: [[__R_2_I]] = phi i64 [ [[ADD28_I]], [[IF_END31_I]] ], [ [[__R_0_I]], [[IF_ELSE17_I]] ] +// CHECK-NEXT: [[COND_I:%.*]] = phi i1 [ true, [[IF_END31_I]] ], [ false, [[IF_ELSE17_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]] ] @@ -128,7 +129,94 @@ // CHECK-LABEL: @test___make_mantissa( // CHECK-NEXT: entry: -// CHECK-NEXT: ret i64 0 +// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[P:%.*]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_I:%.*]] = icmp eq i8 [[TMP0]], 48 +// CHECK-NEXT: br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I32_I:%.*]] +// CHECK: if.then.i: +// CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[P]], i64 1 +// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_I:%.*]] [ +// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_PREHEADER:%.*]] +// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_PREHEADER]] +// CHECK-NEXT: ] +// CHECK: while.cond.i.i.preheader: +// CHECK-NEXT: br label [[WHILE_COND_I_I:%.*]] +// CHECK: while.cond.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I:%.*]], [[CLEANUP_I_I:%.*]] ], [ [[INCDEC_PTR_I]], [[WHILE_COND_I_I_PREHEADER]] ] +// CHECK-NEXT: [[__R_0_I_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], [[CLEANUP_I_I]] ], [ 0, [[WHILE_COND_I_I_PREHEADER]] ] +// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP2]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[WHILE_BODY_I_I:%.*]] +// CHECK: while.body.i.i: +// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48 +// CHECK-NEXT: [[TMP4:%.*]] = icmp ult i8 [[TMP3]], 10 +// CHECK-NEXT: br i1 [[TMP4]], label [[IF_END31_I_I:%.*]], label [[IF_ELSE_I_I:%.*]] +// CHECK: if.else.i.i: +// CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -97 +// CHECK-NEXT: [[TMP6:%.*]] = icmp ult i8 [[TMP5]], 6 +// CHECK-NEXT: br i1 [[TMP6]], label [[IF_END31_I_I]], label [[IF_ELSE17_I_I:%.*]] +// CHECK: if.else17.i.i: +// CHECK-NEXT: [[TMP7:%.*]] = add i8 [[TMP2]], -65 +// CHECK-NEXT: [[TMP8:%.*]] = icmp ult i8 [[TMP7]], 6 +// CHECK-NEXT: br i1 [[TMP8]], label [[IF_END31_I_I]], label [[CLEANUP_I_I]] +// CHECK: if.end31.i.i: +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE17_I_I]] ] +// CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I_I]], 4 +// CHECK-NEXT: [[CONV25_I_I:%.*]] = sext i8 [[TMP2]] to i64 +// CHECK-NEXT: [[ADD26_I_I:%.*]] = add i64 [[MUL24_I_I]], [[DOTSINK]] +// CHECK-NEXT: [[ADD28_I_I:%.*]] = add i64 [[ADD26_I_I]], [[CONV25_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I_I]] +// CHECK: cleanup.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I_I]] = phi ptr [ [[INCDEC_PTR_I_I]], [[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I_I]], [[IF_ELSE17_I_I]] ] +// CHECK-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], [[IF_END31_I_I]] ], [ [[__R_0_I_I]], [[IF_ELSE17_I_I]] ] +// CHECK-NEXT: [[COND_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I]] ], [ false, [[IF_ELSE17_I_I]] ] +// CHECK-NEXT: br i1 [[COND_I_I]], label [[WHILE_COND_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP10]] +// CHECK: while.cond.i17.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I25_I:%.*]], [[CLEANUP_I27_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ] +// CHECK-NEXT: [[__R_0_I15_I:%.*]] = phi i64 [ [[__R_1_I26_I:%.*]], [[CLEANUP_I27_I]] ], [ 0, [[IF_THEN_I]] ] +// CHECK-NEXT: [[TMP9:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I16_I:%.*]] = icmp eq i8 [[TMP9]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I16_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I18_I:%.*]] +// CHECK: while.body.i18.i: +// CHECK-NEXT: [[TMP10:%.*]] = and i8 [[TMP9]], -8 +// CHECK-NEXT: [[TMP11:%.*]] = icmp eq i8 [[TMP10]], 48 +// CHECK-NEXT: br i1 [[TMP11]], label [[IF_THEN_I24_I:%.*]], label [[CLEANUP_I27_I]] +// CHECK: if.then.i24.i: +// CHECK-NEXT: [[MUL_I19_I:%.*]] = shl i64 [[__R_0_I15_I]], 3 +// CHECK-NEXT: [[CONV5_I20_I:%.*]] = sext i8 [[TMP9]] to i64 +// CHECK-NEXT: [[ADD_I21_I:%.*]] = add i64 [[MUL_I19_I]], -48 +// CHECK-NEXT: [[SUB_I22_I:%.*]] = add i64 [[ADD_I21_I]], [[CONV5_I20_I]] +// CHECK-NEXT: [[INCDEC_PTR_I23_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I27_I]] +// CHECK: cleanup.i27.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I25_I]] = phi ptr [ [[INCDEC_PTR_I23_I]], [[IF_THEN_I24_I]] ], [ [[__TAGP_ADDR_0_I14_I]], [[WHILE_BODY_I18_I]] ] +// CHECK-NEXT: [[__R_1_I26_I]] = phi i64 [ [[SUB_I22_I]], [[IF_THEN_I24_I]] ], [ [[__R_0_I15_I]], [[WHILE_BODY_I18_I]] ] +// CHECK-NEXT: br i1 [[TMP11]], label [[WHILE_COND_I17_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP6]] +// CHECK: while.cond.i32.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I29_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I40_I:%.*]], [[CLEANUP_I42_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ] +// CHECK-NEXT: [[__R_0_I30_I:%.*]] = phi i64 [ [[__R_1_I41_I:%.*]], [[CLEANUP_I42_I]] ], [ 0, [[ENTRY]] ] +// CHECK-NEXT: [[TMP12:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I29_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I31_I:%.*]] = icmp eq i8 [[TMP12]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I31_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I33_I:%.*]] +// CHECK: while.body.i33.i: +// CHECK-NEXT: [[TMP13:%.*]] = add i8 [[TMP12]], -48 +// CHECK-NEXT: [[TMP14:%.*]] = icmp ult i8 [[TMP13]], 10 +// CHECK-NEXT: br i1 [[TMP14]], label [[IF_THEN_I39_I:%.*]], label [[CLEANUP_I42_I]] +// CHECK: if.then.i39.i: +// CHECK-NEXT: [[MUL_I34_I:%.*]] = mul i64 [[__R_0_I30_I]], 10 +// CHECK-NEXT: [[CONV5_I35_I:%.*]] = sext i8 [[TMP12]] to i64 +// CHECK-NEXT: [[ADD_I36_I:%.*]] = add i64 [[MUL_I34_I]], -48 +// CHECK-NEXT: [[SUB_I37_I:%.*]] = add i64 [[ADD_I36_I]], [[CONV5_I35_I]] +// CHECK-NEXT: [[INCDEC_PTR_I38_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I29_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I42_I]] +// CHECK: cleanup.i42.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I40_I]] = phi ptr [ [[INCDEC_PTR_I38_I]], [[IF_THEN_I39_I]] ], [ [[__TAGP_ADDR_0_I29_I]], [[WHILE_BODY_I33_I]] ] +// CHECK-NEXT: [[__R_1_I41_I]] = phi i64 [ [[SUB_I37_I]], [[IF_THEN_I39_I]] ], [ [[__R_0_I30_I]], [[WHILE_BODY_I33_I]] ] +// CHECK-NEXT: br i1 [[TMP14]], label [[WHILE_COND_I32_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP9]] +// CHECK: _ZL15__make_mantissaPKc.exit: +// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I27_I]] ], [ [[__R_0_I15_I]], [[WHILE_COND_I17_I]] ], [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I42_I]] ], [ [[__R_0_I30_I]], [[WHILE_COND_I32_I]] ] +// CHECK-NEXT: ret i64 [[RETVAL_0_I]] // extern "C" __device__ uint64_t test___make_mantissa(const char *p) { return __make_mantissa(p); @@ -680,6 +768,7 @@ 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]] @@ -1676,7 +1765,98 @@ // CHECK-LABEL: @test_nanf( // CHECK-NEXT: entry: -// CHECK-NEXT: ret float 0x7FF8000000000000 +// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48 +// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I32_I_I:%.*]] +// CHECK: if.then.i.i: +// CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[TAG]], i64 1 +// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_I_I:%.*]] [ +// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]] +// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_I_PREHEADER]] +// CHECK-NEXT: ] +// CHECK: while.cond.i.i.i.preheader: +// CHECK-NEXT: br label [[WHILE_COND_I_I_I:%.*]] +// CHECK: while.cond.i.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I_I_I_PREHEADER]] ] +// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ] +// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]], label [[WHILE_BODY_I_I_I:%.*]] +// CHECK: while.body.i.i.i: +// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48 +// CHECK-NEXT: [[TMP4:%.*]] = icmp ult i8 [[TMP3]], 10 +// CHECK-NEXT: br i1 [[TMP4]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]] +// CHECK: if.else.i.i.i: +// CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -97 +// CHECK-NEXT: [[TMP6:%.*]] = icmp ult i8 [[TMP5]], 6 +// CHECK-NEXT: br i1 [[TMP6]], label [[IF_END31_I_I_I]], label [[IF_ELSE17_I_I_I:%.*]] +// CHECK: if.else17.i.i.i: +// CHECK-NEXT: [[TMP7:%.*]] = add i8 [[TMP2]], -65 +// CHECK-NEXT: [[TMP8:%.*]] = icmp ult i8 [[TMP7]], 6 +// CHECK-NEXT: br i1 [[TMP8]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]] +// CHECK: if.end31.i.i.i: +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4 +// CHECK-NEXT: [[CONV25_I_I_I:%.*]] = sext i8 [[TMP2]] to i64 +// CHECK-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]] +// CHECK-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I_I_I]] +// CHECK: cleanup.i.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]] +// CHECK: while.cond.i17.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I25_I_I:%.*]], [[CLEANUP_I27_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ] +// CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I26_I_I:%.*]], [[CLEANUP_I27_I_I]] ], [ 0, [[IF_THEN_I_I]] ] +// CHECK-NEXT: [[TMP9:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP9]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I16_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I18_I_I:%.*]] +// CHECK: while.body.i18.i.i: +// CHECK-NEXT: [[TMP10:%.*]] = and i8 [[TMP9]], -8 +// CHECK-NEXT: [[TMP11:%.*]] = icmp eq i8 [[TMP10]], 48 +// CHECK-NEXT: br i1 [[TMP11]], label [[IF_THEN_I24_I_I:%.*]], label [[CLEANUP_I27_I_I]] +// CHECK: if.then.i24.i.i: +// CHECK-NEXT: [[MUL_I19_I_I:%.*]] = shl i64 [[__R_0_I15_I_I]], 3 +// CHECK-NEXT: [[CONV5_I20_I_I:%.*]] = sext i8 [[TMP9]] to i64 +// CHECK-NEXT: [[ADD_I21_I_I:%.*]] = add i64 [[MUL_I19_I_I]], -48 +// CHECK-NEXT: [[SUB_I22_I_I:%.*]] = add i64 [[ADD_I21_I_I]], [[CONV5_I20_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I23_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I27_I_I]] +// CHECK: cleanup.i27.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I25_I_I]] = phi ptr [ [[INCDEC_PTR_I23_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I14_I_I]], [[WHILE_BODY_I18_I_I]] ] +// CHECK-NEXT: [[__R_1_I26_I_I]] = phi i64 [ [[SUB_I22_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_BODY_I18_I_I]] ] +// CHECK-NEXT: br i1 [[TMP11]], label [[WHILE_COND_I17_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP6]] +// CHECK: while.cond.i32.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I29_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I40_I_I:%.*]], [[CLEANUP_I42_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ] +// CHECK-NEXT: [[__R_0_I30_I_I:%.*]] = phi i64 [ [[__R_1_I41_I_I:%.*]], [[CLEANUP_I42_I_I]] ], [ 0, [[ENTRY]] ] +// CHECK-NEXT: [[TMP12:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I29_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I31_I_I:%.*]] = icmp eq i8 [[TMP12]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I31_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I33_I_I:%.*]] +// CHECK: while.body.i33.i.i: +// CHECK-NEXT: [[TMP13:%.*]] = add i8 [[TMP12]], -48 +// CHECK-NEXT: [[TMP14:%.*]] = icmp ult i8 [[TMP13]], 10 +// CHECK-NEXT: br i1 [[TMP14]], label [[IF_THEN_I39_I_I:%.*]], label [[CLEANUP_I42_I_I]] +// CHECK: if.then.i39.i.i: +// CHECK-NEXT: [[MUL_I34_I_I:%.*]] = mul i64 [[__R_0_I30_I_I]], 10 +// CHECK-NEXT: [[CONV5_I35_I_I:%.*]] = sext i8 [[TMP12]] to i64 +// CHECK-NEXT: [[ADD_I36_I_I:%.*]] = add i64 [[MUL_I34_I_I]], -48 +// CHECK-NEXT: [[SUB_I37_I_I:%.*]] = add i64 [[ADD_I36_I_I]], [[CONV5_I35_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I38_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I29_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I42_I_I]] +// CHECK: cleanup.i42.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I40_I_I]] = phi ptr [ [[INCDEC_PTR_I38_I_I]], [[IF_THEN_I39_I_I]] ], [ [[__TAGP_ADDR_0_I29_I_I]], [[WHILE_BODY_I33_I_I]] ] +// CHECK-NEXT: [[__R_1_I41_I_I]] = phi i64 [ [[SUB_I37_I_I]], [[IF_THEN_I39_I_I]] ], [ [[__R_0_I30_I_I]], [[WHILE_BODY_I33_I_I]] ] +// CHECK-NEXT: br i1 [[TMP14]], label [[WHILE_COND_I32_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP9]] +// CHECK: _ZL4nanfPKc.exit: +// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I27_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I42_I_I]] ], [ [[__R_0_I30_I_I]], [[WHILE_COND_I32_I_I]] ] +// CHECK-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32 +// CHECK-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303 +// CHECK-NEXT: [[BF_SET9_I:%.*]] = or i32 [[BF_VALUE_I]], 2143289344 +// CHECK-NEXT: [[TMP15:%.*]] = bitcast i32 [[BF_SET9_I]] to float +// CHECK-NEXT: ret float [[TMP15]] // extern "C" __device__ float test_nanf(const char *tag) { return nanf(tag); @@ -1684,7 +1864,97 @@ // CHECK-LABEL: @test_nan( // CHECK-NEXT: entry: -// CHECK-NEXT: ret double 0x7FF8000000000000 +// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48 +// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I32_I_I:%.*]] +// CHECK: if.then.i.i: +// CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[TAG]], i64 1 +// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_I_I:%.*]] [ +// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]] +// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_I_PREHEADER]] +// CHECK-NEXT: ] +// CHECK: while.cond.i.i.i.preheader: +// CHECK-NEXT: br label [[WHILE_COND_I_I_I:%.*]] +// CHECK: while.cond.i.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I_I_I_PREHEADER]] ] +// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ] +// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT:%.*]], label [[WHILE_BODY_I_I_I:%.*]] +// CHECK: while.body.i.i.i: +// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48 +// CHECK-NEXT: [[TMP4:%.*]] = icmp ult i8 [[TMP3]], 10 +// CHECK-NEXT: br i1 [[TMP4]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]] +// CHECK: if.else.i.i.i: +// CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -97 +// CHECK-NEXT: [[TMP6:%.*]] = icmp ult i8 [[TMP5]], 6 +// CHECK-NEXT: br i1 [[TMP6]], label [[IF_END31_I_I_I]], label [[IF_ELSE17_I_I_I:%.*]] +// CHECK: if.else17.i.i.i: +// CHECK-NEXT: [[TMP7:%.*]] = add i8 [[TMP2]], -65 +// CHECK-NEXT: [[TMP8:%.*]] = icmp ult i8 [[TMP7]], 6 +// CHECK-NEXT: br i1 [[TMP8]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]] +// CHECK: if.end31.i.i.i: +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4 +// CHECK-NEXT: [[CONV25_I_I_I:%.*]] = sext i8 [[TMP2]] to i64 +// CHECK-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]] +// CHECK-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I_I_I]] +// CHECK: cleanup.i.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]] +// CHECK: while.cond.i17.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I25_I_I:%.*]], [[CLEANUP_I27_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ] +// CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I26_I_I:%.*]], [[CLEANUP_I27_I_I]] ], [ 0, [[IF_THEN_I_I]] ] +// CHECK-NEXT: [[TMP9:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP9]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I16_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I18_I_I:%.*]] +// CHECK: while.body.i18.i.i: +// CHECK-NEXT: [[TMP10:%.*]] = and i8 [[TMP9]], -8 +// CHECK-NEXT: [[TMP11:%.*]] = icmp eq i8 [[TMP10]], 48 +// CHECK-NEXT: br i1 [[TMP11]], label [[IF_THEN_I24_I_I:%.*]], label [[CLEANUP_I27_I_I]] +// CHECK: if.then.i24.i.i: +// CHECK-NEXT: [[MUL_I19_I_I:%.*]] = shl i64 [[__R_0_I15_I_I]], 3 +// CHECK-NEXT: [[CONV5_I20_I_I:%.*]] = sext i8 [[TMP9]] to i64 +// CHECK-NEXT: [[ADD_I21_I_I:%.*]] = add i64 [[MUL_I19_I_I]], -48 +// CHECK-NEXT: [[SUB_I22_I_I:%.*]] = add i64 [[ADD_I21_I_I]], [[CONV5_I20_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I23_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I27_I_I]] +// CHECK: cleanup.i27.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I25_I_I]] = phi ptr [ [[INCDEC_PTR_I23_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I14_I_I]], [[WHILE_BODY_I18_I_I]] ] +// CHECK-NEXT: [[__R_1_I26_I_I]] = phi i64 [ [[SUB_I22_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_BODY_I18_I_I]] ] +// CHECK-NEXT: br i1 [[TMP11]], label [[WHILE_COND_I17_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP6]] +// CHECK: while.cond.i32.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I29_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I40_I_I:%.*]], [[CLEANUP_I42_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ] +// CHECK-NEXT: [[__R_0_I30_I_I:%.*]] = phi i64 [ [[__R_1_I41_I_I:%.*]], [[CLEANUP_I42_I_I]] ], [ 0, [[ENTRY]] ] +// CHECK-NEXT: [[TMP12:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I29_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I31_I_I:%.*]] = icmp eq i8 [[TMP12]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I31_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I33_I_I:%.*]] +// CHECK: while.body.i33.i.i: +// CHECK-NEXT: [[TMP13:%.*]] = add i8 [[TMP12]], -48 +// CHECK-NEXT: [[TMP14:%.*]] = icmp ult i8 [[TMP13]], 10 +// CHECK-NEXT: br i1 [[TMP14]], label [[IF_THEN_I39_I_I:%.*]], label [[CLEANUP_I42_I_I]] +// CHECK: if.then.i39.i.i: +// CHECK-NEXT: [[MUL_I34_I_I:%.*]] = mul i64 [[__R_0_I30_I_I]], 10 +// CHECK-NEXT: [[CONV5_I35_I_I:%.*]] = sext i8 [[TMP12]] to i64 +// CHECK-NEXT: [[ADD_I36_I_I:%.*]] = add i64 [[MUL_I34_I_I]], -48 +// CHECK-NEXT: [[SUB_I37_I_I:%.*]] = add i64 [[ADD_I36_I_I]], [[CONV5_I35_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I38_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I29_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I42_I_I]] +// CHECK: cleanup.i42.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I40_I_I]] = phi ptr [ [[INCDEC_PTR_I38_I_I]], [[IF_THEN_I39_I_I]] ], [ [[__TAGP_ADDR_0_I29_I_I]], [[WHILE_BODY_I33_I_I]] ] +// CHECK-NEXT: [[__R_1_I41_I_I]] = phi i64 [ [[SUB_I37_I_I]], [[IF_THEN_I39_I_I]] ], [ [[__R_0_I30_I_I]], [[WHILE_BODY_I33_I_I]] ] +// CHECK-NEXT: br i1 [[TMP14]], label [[WHILE_COND_I32_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP9]] +// CHECK: _ZL3nanPKc.exit: +// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I27_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I42_I_I]] ], [ [[__R_0_I30_I_I]], [[WHILE_COND_I32_I_I]] ] +// CHECK-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247 +// CHECK-NEXT: [[BF_SET9_I:%.*]] = or i64 [[BF_VALUE_I]], 9221120237041090560 +// CHECK-NEXT: [[TMP15:%.*]] = bitcast i64 [[BF_SET9_I]] to double +// CHECK-NEXT: ret double [[TMP15]] // extern "C" __device__ double test_nan(const char *tag) { return nan(tag);