clang/HIP: Fix broken implementations of __make_mantissa* functions

[arsenm]

Summary

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.

[arsenm]

ping

[JonChesterfield]

Patch is obviously fine. Given these are internal functions, perhaps we should annotate them with the non-null attribute and delete the early exit.

[arsenm]

Don't handle null and add nonnull attributes

[arsenm]

ping

[yaxunl]

LGTM. Thanks.

It seems gcc assumes the argument to nan is nonnull (https://godbolt.org/z/xzb8T6Gon), so we can assume that too.

[arsenm]

4086ea331cad827d74542e52a86b7d7933376e7b

clang/lib/Headers/__clang_hip_math.h

 #else
 #define __static_assert_type_size_equal(A,B) \
   static_assert((A) == (B), "")
 
 #endif
 
 __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')
       __r = (__r * 8u) + __tmp - '0';
     else
       return 0;
 
     ++__tagp;
   }
 
   return __r;
 }
 
 __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')
       __r = (__r * 10u) + __tmp - '0';
     else
       return 0;
 
     ++__tagp;
   }
 
   return __r;
 }
 
 __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')
       __r = (__r * 16u) + __tmp - '0';
     else if (__tmp >= 'a' && __tmp <= 'f')
       __r = (__r * 16u) + __tmp - 'a' + 10;
     else if (__tmp >= 'A' && __tmp <= 'F')
       __r = (__r * 16u) + __tmp - 'A' + 10;

clang/test/Headers/__clang_hip_math.hip

 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -ffinite-math-only -o - \
 // RUN:   -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,FINITEONLY %s
 
 #define BOOL_TYPE int
 typedef unsigned long long uint64_t;
 
 // CHECK-LABEL: @test___make_mantissa_base8(
 // CHECK-NEXT:  entry:
-// CHECK-NEXT:    br label [[WHILE_COND_I:%.*]]
+// CHECK-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);
 }
 
 // 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);
 }
 
 // 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);
 }
 
 // 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);
 }
 
 // CHECK-LABEL: @test_abs(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.abs.i32(i32 [[X:%.*]], i1 true)
 // CHECK-NEXT:    ret i32 [[TMP0]]
 //
 extern "C" __device__ int test_abs(int x) {
   return abs(x);
 }
 
 // CHECK-LABEL: @test_labs(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
 // CHECK-NEXT:    ret i64 [[TMP0]]
 //
 extern "C" __device__ long test_labs(long x) {
   return labs(x);
 }
 
-//
 // CHECK-LABEL: @test_llabs(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
 // CHECK-NEXT:    ret i64 [[TMP0]]
 //
 extern "C" __device__ long long test_llabs(long x) {
   return llabs(x);
 }

 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
 extern "C" __device__ double test_modf(double x, double* y) {
   return modf(x, y);
 }
 
 // CHECK-LABEL: @test_nanf(
 // CHECK-NEXT:  entry:
-// CHECK-NEXT:    ret float 0x7FF8000000000000
+// 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);
 }
 
 // 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);
 }
 
 // CHECK-LABEL: @test_nanf_emptystr(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret float 0x7FF8000000000000