arsenm created this revision.
arsenm added reviewers: yaxunl, JonChesterfield, saiislam, scchan, AlexVlx,
b-sumner.
Herald added a project: All.
arsenm requested review of this revision.
Herald added a subscriber: wdng.
The optimizer was folding the entire function to return 0. This
meant to be checking the character content of the pointer is the
string terminator, not null.
Not sure if the null checks are necessary. I'm not sure what
the spec is for these, or why these would even be exposed.
https://reviews.llvm.org/D138392
Files:
clang/lib/Headers/__clang_hip_math.h
clang/test/Headers/__clang_hip_math.hip
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);
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')
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
