arsenm updated this revision to Diff 482839.
arsenm added a comment.
Don't handle null and add nonnull attributes
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D138392/new/
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
@@ -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);
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 {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
