Index: clang/lib/Headers/__clang_hip_math.h =================================================================== --- clang/lib/Headers/__clang_hip_math.h +++ clang/lib/Headers/__clang_hip_math.h @@ -71,8 +71,11 @@ __DEVICE__ uint64_t __make_mantissa_base8(const char *__tagp) { + if (!__tagp) + return 0; + uint64_t __r = 0; - while (__tagp) { + while (*__tagp != '\0') { char __tmp = *__tagp; if (__tmp >= '0' && __tmp <= '7') @@ -88,8 +91,11 @@ __DEVICE__ uint64_t __make_mantissa_base10(const char *__tagp) { + if (!__tagp) + return 0; + uint64_t __r = 0; - while (__tagp) { + while (*__tagp != '\0') { char __tmp = *__tagp; if (__tmp >= '0' && __tmp <= '9') @@ -105,8 +111,11 @@ __DEVICE__ uint64_t __make_mantissa_base16(const char *__tagp) { + if (!__tagp) + return 0; + uint64_t __r = 0; - while (__tagp) { + while (*__tagp != '\0') { char __tmp = *__tagp; if (__tmp >= '0' && __tmp <= '9') Index: clang/test/Headers/__clang_hip_math.hip =================================================================== --- clang/test/Headers/__clang_hip_math.hip +++ clang/test/Headers/__clang_hip_math.hip @@ -22,31 +22,32 @@ // CHECK-LABEL: @test___make_mantissa_base8( // CHECK-NEXT: entry: -// CHECK-NEXT: br label [[WHILE_COND_I:%.*]] +// CHECK-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq ptr [[P:%.*]], null +// CHECK-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT:%.*]], 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: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ] +// CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ [[__R_1_I:%.*]], [[CLEANUP_I]] ], [ 0, [[ENTRY]] ] // 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: br i1 [[TMP2]], label [[IF_THEN5_I:%.*]], label [[CLEANUP_I]] +// CHECK: if.then5.i: // CHECK-NEXT: [[MUL_I:%.*]] = shl i64 [[__R_0_I]], 3 -// CHECK-NEXT: [[CONV3_I:%.*]] = sext i8 [[TMP0]] to i64 +// CHECK-NEXT: [[CONV6_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]], [[CONV6_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: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_THEN5_I]] ], [ [[__TAGP_ADDR_0_I]], [[WHILE_BODY_I]] ] +// CHECK-NEXT: [[__R_1_I]] = phi i64 [ [[SUB_I]], [[IF_THEN5_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]] +// CHECK-NEXT: [[RETVAL_3_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ] +// CHECK-NEXT: ret i64 [[RETVAL_3_I]] // extern "C" __device__ uint64_t test___make_mantissa_base8(const char *p) { return __make_mantissa_base8(p); @@ -54,31 +55,32 @@ // CHECK-LABEL: @test___make_mantissa_base10( // CHECK-NEXT: entry: -// CHECK-NEXT: br label [[WHILE_COND_I:%.*]] +// CHECK-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq ptr [[P:%.*]], null +// CHECK-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT:%.*]], 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: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ] +// CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ [[__R_1_I:%.*]], [[CLEANUP_I]] ], [ 0, [[ENTRY]] ] // 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: br i1 [[TMP2]], label [[IF_THEN5_I:%.*]], label [[CLEANUP_I]] +// CHECK: if.then5.i: // CHECK-NEXT: [[MUL_I:%.*]] = mul i64 [[__R_0_I]], 10 -// CHECK-NEXT: [[CONV3_I:%.*]] = sext i8 [[TMP0]] to i64 +// CHECK-NEXT: [[CONV6_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]], [[CONV6_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: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_THEN5_I]] ], [ [[__TAGP_ADDR_0_I]], [[WHILE_BODY_I]] ] +// CHECK-NEXT: [[__R_1_I]] = phi i64 [ [[SUB_I]], [[IF_THEN5_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]] +// CHECK-NEXT: [[RETVAL_3_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ] +// CHECK-NEXT: ret i64 [[RETVAL_3_I]] // extern "C" __device__ uint64_t test___make_mantissa_base10(const char *p) { return __make_mantissa_base10(p); @@ -86,41 +88,42 @@ // CHECK-LABEL: @test___make_mantissa_base16( // CHECK-NEXT: entry: -// CHECK-NEXT: br label [[WHILE_COND_I:%.*]] +// CHECK-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq ptr [[P:%.*]], null +// CHECK-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT:%.*]], 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: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ] +// CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ [[__R_2_I:%.*]], [[CLEANUP_I]] ], [ 0, [[ENTRY]] ] // 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_END33_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_END33_I]], label [[IF_ELSE18_I:%.*]] +// CHECK: if.else18.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_END33_I]], label [[CLEANUP_I]] +// CHECK: if.end33.i: +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I]] ], [ -87, [[IF_ELSE_I]] ], [ -55, [[IF_ELSE18_I]] ] +// CHECK-NEXT: [[MUL25_I:%.*]] = shl i64 [[__R_0_I]], 4 +// CHECK-NEXT: [[CONV26_I:%.*]] = sext i8 [[TMP0]] to i64 +// CHECK-NEXT: [[ADD27_I:%.*]] = add i64 [[MUL25_I]], [[DOTSINK]] +// CHECK-NEXT: [[ADD29_I:%.*]] = add i64 [[ADD27_I]], [[CONV26_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_END33_I]] ], [ [[__TAGP_ADDR_0_I]], [[IF_ELSE18_I]] ] +// CHECK-NEXT: [[__R_2_I]] = phi i64 [ [[ADD29_I]], [[IF_END33_I]] ], [ [[__R_0_I]], [[IF_ELSE18_I]] ] +// CHECK-NEXT: [[COND_I:%.*]] = phi i1 [ true, [[IF_END33_I]] ], [ false, [[IF_ELSE18_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]] +// CHECK-NEXT: [[RETVAL_3_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ] +// CHECK-NEXT: ret i64 [[RETVAL_3_I]] // extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) { return __make_mantissa_base16(p); @@ -128,7 +131,97 @@ // CHECK-LABEL: @test___make_mantissa( // CHECK-NEXT: entry: -// CHECK-NEXT: ret i64 0 +// CHECK-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq ptr [[P:%.*]], null +// CHECK-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[IF_END_I:%.*]] +// CHECK: if.end.i: +// 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_THEN1_I:%.*]], label [[WHILE_COND_I35_I:%.*]] +// CHECK: if.then1.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_I20_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_END33_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_END33_I_I]], label [[IF_ELSE18_I_I:%.*]] +// CHECK: if.else18.i.i: +// CHECK-NEXT: [[TMP7:%.*]] = add i8 [[TMP2]], -65 +// CHECK-NEXT: [[TMP8:%.*]] = icmp ult i8 [[TMP7]], 6 +// CHECK-NEXT: br i1 [[TMP8]], label [[IF_END33_I_I]], label [[CLEANUP_I_I]] +// CHECK: if.end33.i.i: +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE18_I_I]] ] +// CHECK-NEXT: [[MUL25_I_I:%.*]] = shl i64 [[__R_0_I_I]], 4 +// CHECK-NEXT: [[CONV26_I_I:%.*]] = sext i8 [[TMP2]] to i64 +// CHECK-NEXT: [[ADD27_I_I:%.*]] = add i64 [[MUL25_I_I]], [[DOTSINK]] +// CHECK-NEXT: [[ADD29_I_I:%.*]] = add i64 [[ADD27_I_I]], [[CONV26_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_END33_I_I]] ], [ [[__TAGP_ADDR_0_I_I]], [[IF_ELSE18_I_I]] ] +// CHECK-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD29_I_I]], [[IF_END33_I_I]] ], [ [[__R_0_I_I]], [[IF_ELSE18_I_I]] ] +// CHECK-NEXT: [[COND_I_I:%.*]] = phi i1 [ true, [[IF_END33_I_I]] ], [ false, [[IF_ELSE18_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.i20.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I17_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I28_I:%.*]], [[CLEANUP_I30_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN1_I]] ] +// CHECK-NEXT: [[__R_0_I18_I:%.*]] = phi i64 [ [[__R_1_I29_I:%.*]], [[CLEANUP_I30_I]] ], [ 0, [[IF_THEN1_I]] ] +// CHECK-NEXT: [[TMP9:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I17_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I19_I:%.*]] = icmp eq i8 [[TMP9]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I19_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I21_I:%.*]] +// CHECK: while.body.i21.i: +// CHECK-NEXT: [[TMP10:%.*]] = and i8 [[TMP9]], -8 +// CHECK-NEXT: [[TMP11:%.*]] = icmp eq i8 [[TMP10]], 48 +// CHECK-NEXT: br i1 [[TMP11]], label [[IF_THEN5_I27_I:%.*]], label [[CLEANUP_I30_I]] +// CHECK: if.then5.i27.i: +// CHECK-NEXT: [[MUL_I22_I:%.*]] = shl i64 [[__R_0_I18_I]], 3 +// CHECK-NEXT: [[CONV6_I23_I:%.*]] = sext i8 [[TMP9]] to i64 +// CHECK-NEXT: [[ADD_I24_I:%.*]] = add i64 [[MUL_I22_I]], -48 +// CHECK-NEXT: [[SUB_I25_I:%.*]] = add i64 [[ADD_I24_I]], [[CONV6_I23_I]] +// CHECK-NEXT: [[INCDEC_PTR_I26_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I17_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I30_I]] +// CHECK: cleanup.i30.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I28_I]] = phi ptr [ [[INCDEC_PTR_I26_I]], [[IF_THEN5_I27_I]] ], [ [[__TAGP_ADDR_0_I17_I]], [[WHILE_BODY_I21_I]] ] +// CHECK-NEXT: [[__R_1_I29_I]] = phi i64 [ [[SUB_I25_I]], [[IF_THEN5_I27_I]] ], [ [[__R_0_I18_I]], [[WHILE_BODY_I21_I]] ] +// CHECK-NEXT: br i1 [[TMP11]], label [[WHILE_COND_I20_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP6]] +// CHECK: while.cond.i35.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I32_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I:%.*]], [[CLEANUP_I45_I:%.*]] ], [ [[P]], [[IF_END_I]] ] +// CHECK-NEXT: [[__R_0_I33_I:%.*]] = phi i64 [ [[__R_1_I44_I:%.*]], [[CLEANUP_I45_I]] ], [ 0, [[IF_END_I]] ] +// CHECK-NEXT: [[TMP12:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I32_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I34_I:%.*]] = icmp eq i8 [[TMP12]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I34_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I36_I:%.*]] +// CHECK: while.body.i36.i: +// CHECK-NEXT: [[TMP13:%.*]] = add i8 [[TMP12]], -48 +// CHECK-NEXT: [[TMP14:%.*]] = icmp ult i8 [[TMP13]], 10 +// CHECK-NEXT: br i1 [[TMP14]], label [[IF_THEN5_I42_I:%.*]], label [[CLEANUP_I45_I]] +// CHECK: if.then5.i42.i: +// CHECK-NEXT: [[MUL_I37_I:%.*]] = mul i64 [[__R_0_I33_I]], 10 +// CHECK-NEXT: [[CONV6_I38_I:%.*]] = sext i8 [[TMP12]] to i64 +// CHECK-NEXT: [[ADD_I39_I:%.*]] = add i64 [[MUL_I37_I]], -48 +// CHECK-NEXT: [[SUB_I40_I:%.*]] = add i64 [[ADD_I39_I]], [[CONV6_I38_I]] +// CHECK-NEXT: [[INCDEC_PTR_I41_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I32_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I45_I]] +// CHECK: cleanup.i45.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I43_I]] = phi ptr [ [[INCDEC_PTR_I41_I]], [[IF_THEN5_I42_I]] ], [ [[__TAGP_ADDR_0_I32_I]], [[WHILE_BODY_I36_I]] ] +// CHECK-NEXT: [[__R_1_I44_I]] = phi i64 [ [[SUB_I40_I]], [[IF_THEN5_I42_I]] ], [ [[__R_0_I33_I]], [[WHILE_BODY_I36_I]] ] +// CHECK-NEXT: br i1 [[TMP14]], label [[WHILE_COND_I35_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP9]] +// CHECK: _ZL15__make_mantissaPKc.exit: +// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ 0, [[CLEANUP_I30_I]] ], [ [[__R_0_I18_I]], [[WHILE_COND_I20_I]] ], [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I45_I]] ], [ [[__R_0_I33_I]], [[WHILE_COND_I35_I]] ] +// CHECK-NEXT: ret i64 [[RETVAL_0_I]] // extern "C" __device__ uint64_t test___make_mantissa(const char *p) { return __make_mantissa(p); @@ -152,7 +245,6 @@ return labs(x); } -// // CHECK-LABEL: @test_llabs( // CHECK-NEXT: entry: // CHECK-NEXT: [[TMP0:%.*]] = tail call i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true) @@ -1629,7 +1721,101 @@ // CHECK-LABEL: @test_nanf( // CHECK-NEXT: entry: -// CHECK-NEXT: ret float 0x7FF8000000000000 +// CHECK-NEXT: [[TOBOOL_NOT_I_I:%.*]] = icmp eq ptr [[TAG:%.*]], null +// CHECK-NEXT: br i1 [[TOBOOL_NOT_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]], label [[IF_END_I_I:%.*]] +// CHECK: if.end.i.i: +// 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_THEN1_I_I:%.*]], label [[WHILE_COND_I35_I_I:%.*]] +// CHECK: if.then1.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_I20_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_END33_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_END33_I_I_I]], label [[IF_ELSE18_I_I_I:%.*]] +// CHECK: if.else18.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_END33_I_I_I]], label [[CLEANUP_I_I_I]] +// CHECK: if.end33.i.i.i: +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE18_I_I_I]] ] +// CHECK-NEXT: [[MUL25_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4 +// CHECK-NEXT: [[CONV26_I_I_I:%.*]] = sext i8 [[TMP2]] to i64 +// CHECK-NEXT: [[ADD27_I_I_I:%.*]] = add i64 [[MUL25_I_I_I]], [[DOTSINK]] +// CHECK-NEXT: [[ADD29_I_I_I:%.*]] = add i64 [[ADD27_I_I_I]], [[CONV26_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_END33_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE18_I_I_I]] ] +// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD29_I_I_I]], [[IF_END33_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE18_I_I_I]] ] +// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END33_I_I_I]] ], [ false, [[IF_ELSE18_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.i20.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I17_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I28_I_I:%.*]], [[CLEANUP_I30_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN1_I_I]] ] +// CHECK-NEXT: [[__R_0_I18_I_I:%.*]] = phi i64 [ [[__R_1_I29_I_I:%.*]], [[CLEANUP_I30_I_I]] ], [ 0, [[IF_THEN1_I_I]] ] +// CHECK-NEXT: [[TMP9:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I17_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I19_I_I:%.*]] = icmp eq i8 [[TMP9]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I19_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I21_I_I:%.*]] +// CHECK: while.body.i21.i.i: +// CHECK-NEXT: [[TMP10:%.*]] = and i8 [[TMP9]], -8 +// CHECK-NEXT: [[TMP11:%.*]] = icmp eq i8 [[TMP10]], 48 +// CHECK-NEXT: br i1 [[TMP11]], label [[IF_THEN5_I27_I_I:%.*]], label [[CLEANUP_I30_I_I]] +// CHECK: if.then5.i27.i.i: +// CHECK-NEXT: [[MUL_I22_I_I:%.*]] = shl i64 [[__R_0_I18_I_I]], 3 +// CHECK-NEXT: [[CONV6_I23_I_I:%.*]] = sext i8 [[TMP9]] to i64 +// CHECK-NEXT: [[ADD_I24_I_I:%.*]] = add i64 [[MUL_I22_I_I]], -48 +// CHECK-NEXT: [[SUB_I25_I_I:%.*]] = add i64 [[ADD_I24_I_I]], [[CONV6_I23_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I26_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I17_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I30_I_I]] +// CHECK: cleanup.i30.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I28_I_I]] = phi ptr [ [[INCDEC_PTR_I26_I_I]], [[IF_THEN5_I27_I_I]] ], [ [[__TAGP_ADDR_0_I17_I_I]], [[WHILE_BODY_I21_I_I]] ] +// CHECK-NEXT: [[__R_1_I29_I_I]] = phi i64 [ [[SUB_I25_I_I]], [[IF_THEN5_I27_I_I]] ], [ [[__R_0_I18_I_I]], [[WHILE_BODY_I21_I_I]] ] +// CHECK-NEXT: br i1 [[TMP11]], label [[WHILE_COND_I20_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP6]] +// CHECK: while.cond.i35.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I32_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I_I:%.*]], [[CLEANUP_I45_I_I:%.*]] ], [ [[TAG]], [[IF_END_I_I]] ] +// CHECK-NEXT: [[__R_0_I33_I_I:%.*]] = phi i64 [ [[__R_1_I44_I_I:%.*]], [[CLEANUP_I45_I_I]] ], [ 0, [[IF_END_I_I]] ] +// CHECK-NEXT: [[TMP12:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I32_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I34_I_I:%.*]] = icmp eq i8 [[TMP12]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I34_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I36_I_I:%.*]] +// CHECK: while.body.i36.i.i: +// CHECK-NEXT: [[TMP13:%.*]] = add i8 [[TMP12]], -48 +// CHECK-NEXT: [[TMP14:%.*]] = icmp ult i8 [[TMP13]], 10 +// CHECK-NEXT: br i1 [[TMP14]], label [[IF_THEN5_I42_I_I:%.*]], label [[CLEANUP_I45_I_I]] +// CHECK: if.then5.i42.i.i: +// CHECK-NEXT: [[MUL_I37_I_I:%.*]] = mul i64 [[__R_0_I33_I_I]], 10 +// CHECK-NEXT: [[CONV6_I38_I_I:%.*]] = sext i8 [[TMP12]] to i64 +// CHECK-NEXT: [[ADD_I39_I_I:%.*]] = add i64 [[MUL_I37_I_I]], -48 +// CHECK-NEXT: [[SUB_I40_I_I:%.*]] = add i64 [[ADD_I39_I_I]], [[CONV6_I38_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I41_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I32_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I45_I_I]] +// CHECK: cleanup.i45.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I43_I_I]] = phi ptr [ [[INCDEC_PTR_I41_I_I]], [[IF_THEN5_I42_I_I]] ], [ [[__TAGP_ADDR_0_I32_I_I]], [[WHILE_BODY_I36_I_I]] ] +// CHECK-NEXT: [[__R_1_I44_I_I]] = phi i64 [ [[SUB_I40_I_I]], [[IF_THEN5_I42_I_I]] ], [ [[__R_0_I33_I_I]], [[WHILE_BODY_I36_I_I]] ] +// CHECK-NEXT: br i1 [[TMP14]], label [[WHILE_COND_I35_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP9]] +// CHECK: _ZL4nanfPKc.exit: +// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ 0, [[CLEANUP_I30_I_I]] ], [ [[__R_0_I18_I_I]], [[WHILE_COND_I20_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I45_I_I]] ], [ [[__R_0_I33_I_I]], [[WHILE_COND_I35_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); @@ -1637,7 +1823,100 @@ // CHECK-LABEL: @test_nan( // CHECK-NEXT: entry: -// CHECK-NEXT: ret double 0x7FF8000000000000 +// CHECK-NEXT: [[TOBOOL_NOT_I_I:%.*]] = icmp eq ptr [[TAG:%.*]], null +// CHECK-NEXT: br i1 [[TOBOOL_NOT_I_I]], label [[_ZL3NANPKC_EXIT:%.*]], label [[IF_END_I_I:%.*]] +// CHECK: if.end.i.i: +// 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_THEN1_I_I:%.*]], label [[WHILE_COND_I35_I_I:%.*]] +// CHECK: if.then1.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_I20_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_END33_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_END33_I_I_I]], label [[IF_ELSE18_I_I_I:%.*]] +// CHECK: if.else18.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_END33_I_I_I]], label [[CLEANUP_I_I_I]] +// CHECK: if.end33.i.i.i: +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE18_I_I_I]] ] +// CHECK-NEXT: [[MUL25_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4 +// CHECK-NEXT: [[CONV26_I_I_I:%.*]] = sext i8 [[TMP2]] to i64 +// CHECK-NEXT: [[ADD27_I_I_I:%.*]] = add i64 [[MUL25_I_I_I]], [[DOTSINK]] +// CHECK-NEXT: [[ADD29_I_I_I:%.*]] = add i64 [[ADD27_I_I_I]], [[CONV26_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_END33_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE18_I_I_I]] ] +// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD29_I_I_I]], [[IF_END33_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE18_I_I_I]] ] +// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END33_I_I_I]] ], [ false, [[IF_ELSE18_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.i20.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I17_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I28_I_I:%.*]], [[CLEANUP_I30_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN1_I_I]] ] +// CHECK-NEXT: [[__R_0_I18_I_I:%.*]] = phi i64 [ [[__R_1_I29_I_I:%.*]], [[CLEANUP_I30_I_I]] ], [ 0, [[IF_THEN1_I_I]] ] +// CHECK-NEXT: [[TMP9:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I17_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I19_I_I:%.*]] = icmp eq i8 [[TMP9]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I19_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I21_I_I:%.*]] +// CHECK: while.body.i21.i.i: +// CHECK-NEXT: [[TMP10:%.*]] = and i8 [[TMP9]], -8 +// CHECK-NEXT: [[TMP11:%.*]] = icmp eq i8 [[TMP10]], 48 +// CHECK-NEXT: br i1 [[TMP11]], label [[IF_THEN5_I27_I_I:%.*]], label [[CLEANUP_I30_I_I]] +// CHECK: if.then5.i27.i.i: +// CHECK-NEXT: [[MUL_I22_I_I:%.*]] = shl i64 [[__R_0_I18_I_I]], 3 +// CHECK-NEXT: [[CONV6_I23_I_I:%.*]] = sext i8 [[TMP9]] to i64 +// CHECK-NEXT: [[ADD_I24_I_I:%.*]] = add i64 [[MUL_I22_I_I]], -48 +// CHECK-NEXT: [[SUB_I25_I_I:%.*]] = add i64 [[ADD_I24_I_I]], [[CONV6_I23_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I26_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I17_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I30_I_I]] +// CHECK: cleanup.i30.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I28_I_I]] = phi ptr [ [[INCDEC_PTR_I26_I_I]], [[IF_THEN5_I27_I_I]] ], [ [[__TAGP_ADDR_0_I17_I_I]], [[WHILE_BODY_I21_I_I]] ] +// CHECK-NEXT: [[__R_1_I29_I_I]] = phi i64 [ [[SUB_I25_I_I]], [[IF_THEN5_I27_I_I]] ], [ [[__R_0_I18_I_I]], [[WHILE_BODY_I21_I_I]] ] +// CHECK-NEXT: br i1 [[TMP11]], label [[WHILE_COND_I20_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP6]] +// CHECK: while.cond.i35.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I32_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I_I:%.*]], [[CLEANUP_I45_I_I:%.*]] ], [ [[TAG]], [[IF_END_I_I]] ] +// CHECK-NEXT: [[__R_0_I33_I_I:%.*]] = phi i64 [ [[__R_1_I44_I_I:%.*]], [[CLEANUP_I45_I_I]] ], [ 0, [[IF_END_I_I]] ] +// CHECK-NEXT: [[TMP12:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I32_I_I]], align 1, !tbaa [[TBAA3]] +// CHECK-NEXT: [[CMP_NOT_I34_I_I:%.*]] = icmp eq i8 [[TMP12]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I34_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I36_I_I:%.*]] +// CHECK: while.body.i36.i.i: +// CHECK-NEXT: [[TMP13:%.*]] = add i8 [[TMP12]], -48 +// CHECK-NEXT: [[TMP14:%.*]] = icmp ult i8 [[TMP13]], 10 +// CHECK-NEXT: br i1 [[TMP14]], label [[IF_THEN5_I42_I_I:%.*]], label [[CLEANUP_I45_I_I]] +// CHECK: if.then5.i42.i.i: +// CHECK-NEXT: [[MUL_I37_I_I:%.*]] = mul i64 [[__R_0_I33_I_I]], 10 +// CHECK-NEXT: [[CONV6_I38_I_I:%.*]] = sext i8 [[TMP12]] to i64 +// CHECK-NEXT: [[ADD_I39_I_I:%.*]] = add i64 [[MUL_I37_I_I]], -48 +// CHECK-NEXT: [[SUB_I40_I_I:%.*]] = add i64 [[ADD_I39_I_I]], [[CONV6_I38_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I41_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I32_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I45_I_I]] +// CHECK: cleanup.i45.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I43_I_I]] = phi ptr [ [[INCDEC_PTR_I41_I_I]], [[IF_THEN5_I42_I_I]] ], [ [[__TAGP_ADDR_0_I32_I_I]], [[WHILE_BODY_I36_I_I]] ] +// CHECK-NEXT: [[__R_1_I44_I_I]] = phi i64 [ [[SUB_I40_I_I]], [[IF_THEN5_I42_I_I]] ], [ [[__R_0_I33_I_I]], [[WHILE_BODY_I36_I_I]] ] +// CHECK-NEXT: br i1 [[TMP14]], label [[WHILE_COND_I35_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP9]] +// CHECK: _ZL3nanPKc.exit: +// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[ENTRY:%.*]] ], [ 0, [[CLEANUP_I30_I_I]] ], [ [[__R_0_I18_I_I]], [[WHILE_COND_I20_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I45_I_I]] ], [ [[__R_0_I33_I_I]], [[WHILE_COND_I35_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);