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) {
+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')
       __r = (__r * 8u) + __tmp - '0';
     else
       return 0;
 
     ++__tagp;
   }
 
   return __r;
 }
 
 __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')
       __r = (__r * 10u) + __tmp - '0';
     else
       return 0;
 
     ++__tagp;
   }
 
   return __r;
 }
 
 __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')
       __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;
     else
       return 0;
 
     ++__tagp;
   }
 
   return __r;
 }
 
 __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;
 
     if (*__tagp == 'x' || *__tagp == 'X')
       return __make_mantissa_base16(__tagp);
     else
       return __make_mantissa_base8(__tagp);
   }

 #endif
   float __r =
       __ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
   *__iptr = __tmp;
   return __r;
 }
 
 __DEVICE__
-float nanf(const char *__tagp) {
+float nanf(const char *__tagp __attribute__((nonnull))) {
   union {
     float val;
     struct ieee_float {
       unsigned int mantissa : 22;
       unsigned int quiet : 1;
       unsigned int exponent : 8;
       unsigned int sign : 1;
     } bits;

clang/test/Headers/__clang_hip_math.hip

 typedef unsigned long long uint64_t;
 
 // CHECK-LABEL: @test___make_mantissa_base8(
 // 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_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:
 // 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:    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]]
 //
 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:       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:
 // 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:    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]]
 //
 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:       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]] ]
 // CHECK-NEXT:    ret i64 [[RETVAL_2_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:    [[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);
 }
 
 // CHECK-LABEL: @test_abs(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.abs.i32(i32 [[X:%.*]], i1 true)

 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
 extern "C" __device__ float test_exp2f(float x) {
   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]]
 // DEFAULT-NEXT:    ret double [[CALL_I]]
 //
 // FINITEONLY-LABEL: @test_exp2(
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR15]]

 // 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:    [[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);
 }
 
 // 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);
 }
 
 // CHECK-LABEL: @test_nanf_emptystr(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    ret float 0x7FF8000000000000