clang/test/CodeGenHIP/printf_nonhostcall.cpp

   1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 2
   2 // REQUIRES: amdgpu-registered-target
   3 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -mprintf-kind=buffered -fcuda-is-device \
   4 // RUN:   -o - %s | FileCheck --enable-var-scope %s
   5 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -ffp-exception-behavior=strict -mprintf-kind=buffered -fcuda-is-device \
   6 // RUN:   -o - %s | FileCheck --enable-var-scope --check-prefix=CHECK_CONSTRAINED %s
   7
   8 #define __device__ __attribute__((device))
   9 #define __shared__ __attribute__((shared))
  10 #define __constant__ __attribute__((constant))
  11
  12 extern "C" __device__ int printf(const char *format, ...);
  13
  14 // CHECK-LABEL: define dso_local noundef i32 @_Z4foo1v
  15 // CHECK-NEXT:  entry:
  16 // CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
  17 // CHECK-NEXT:    [[S:%.*]] = alloca ptr, align 8, addrspace(5)
  18 // CHECK-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
  19 // CHECK-NEXT:    [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr
  20 // CHECK-NEXT:    store ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr [[S_ASCAST]], align 8
  21 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
  22 // CHECK-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
  23 // CHECK-NEXT:    [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null
  24 // CHECK-NEXT:    br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
  25 // CHECK:       strlen.while:
  26 // CHECK-NEXT:    [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]
  27 // CHECK-NEXT:    [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1
  28 // CHECK-NEXT:    [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1
  29 // CHECK-NEXT:    [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0
  30 // CHECK-NEXT:    br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
  31 // CHECK:       strlen.while.done:
  32 // CHECK-NEXT:    [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64
  33 // CHECK-NEXT:    [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64
  34 // CHECK-NEXT:    [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]]
  35 // CHECK-NEXT:    [[TMP10:%.*]] = add i64 [[TMP9]], 1
  36 // CHECK-NEXT:    br label [[STRLEN_JOIN]]
  37 // CHECK:       strlen.join:
  38 // CHECK-NEXT:    [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
  39 // CHECK-NEXT:    [[TMP12:%.*]] = add i64 [[TMP11]], 7
  40 // CHECK-NEXT:    [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288
  41 // CHECK-NEXT:    [[TMP14:%.*]] = add i64 [[TMP13]], 52
  42 // CHECK-NEXT:    [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32
  43 // CHECK-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])
  44 // CHECK-NEXT:    [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
  45 // CHECK-NEXT:    br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
  46 // CHECK:       end.block:
  47 // CHECK-NEXT:    [[TMP17:%.*]] = xor i1 [[TMP16]], true
  48 // CHECK-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32
  49 // CHECK-NEXT:    ret i32 [[PRINTF_RESULT]]
  50 // CHECK:       argpush.block:
  51 // CHECK-NEXT:    [[TMP18:%.*]] = shl i32 [[TMP15]], 2
  52 // CHECK-NEXT:    [[TMP19:%.*]] = or i32 [[TMP18]], 2
  53 // CHECK-NEXT:    store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
  54 // CHECK-NEXT:    [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
  55 // CHECK-NEXT:    store i64 1107004088646384690, ptr addrspace(1) [[TMP20]], align 8
  56 // CHECK-NEXT:    [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8
  57 // CHECK-NEXT:    store i64 8, ptr addrspace(1) [[TMP21]], align 8
  58 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i32 8
  59 // CHECK-NEXT:    store double 3.141590e+00, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
  60 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
  61 // CHECK-NEXT:    store i64 8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8
  62 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8
  63 // CHECK-NEXT:    store i64 4, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8
  64 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8
  65 // CHECK-NEXT:    call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[PRINTBUFFNEXTPTR3]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false)
  66 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i64 [[TMP13]]
  67 // CHECK-NEXT:    store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8
  68 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8
  69 // CHECK-NEXT:    br label [[END_BLOCK]]
  70 //
  71 // CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo1v
  72 // CHECK_CONSTRAINED-NEXT:  entry:
  73 // CHECK_CONSTRAINED-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
  74 // CHECK_CONSTRAINED-NEXT:    [[S:%.*]] = alloca ptr, align 8, addrspace(5)
  75 // CHECK_CONSTRAINED-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
  76 // CHECK_CONSTRAINED-NEXT:    [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr
  77 // CHECK_CONSTRAINED-NEXT:    store ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr [[S_ASCAST]], align 8
  78 // CHECK_CONSTRAINED-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
  79 // CHECK_CONSTRAINED-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
  80 // CHECK_CONSTRAINED-NEXT:    [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null
  81 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
  82 // CHECK_CONSTRAINED:       strlen.while:
  83 // CHECK_CONSTRAINED-NEXT:    [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]
  84 // CHECK_CONSTRAINED-NEXT:    [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1
  85 // CHECK_CONSTRAINED-NEXT:    [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1
  86 // CHECK_CONSTRAINED-NEXT:    [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0
  87 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
  88 // CHECK_CONSTRAINED:       strlen.while.done:
  89 // CHECK_CONSTRAINED-NEXT:    [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64
  90 // CHECK_CONSTRAINED-NEXT:    [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64
  91 // CHECK_CONSTRAINED-NEXT:    [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]]
  92 // CHECK_CONSTRAINED-NEXT:    [[TMP10:%.*]] = add i64 [[TMP9]], 1
  93 // CHECK_CONSTRAINED-NEXT:    br label [[STRLEN_JOIN]]
  94 // CHECK_CONSTRAINED:       strlen.join:
  95 // CHECK_CONSTRAINED-NEXT:    [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
  96 // CHECK_CONSTRAINED-NEXT:    [[TMP12:%.*]] = add i64 [[TMP11]], 7
  97 // CHECK_CONSTRAINED-NEXT:    [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288
  98 // CHECK_CONSTRAINED-NEXT:    [[TMP14:%.*]] = add i64 [[TMP13]], 52
  99 // CHECK_CONSTRAINED-NEXT:    [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32
 100 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])
 101 // CHECK_CONSTRAINED-NEXT:    [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
 102 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
 103 // CHECK_CONSTRAINED:       end.block:
 104 // CHECK_CONSTRAINED-NEXT:    [[TMP17:%.*]] = xor i1 [[TMP16]], true
 105 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32
 106 // CHECK_CONSTRAINED-NEXT:    ret i32 [[PRINTF_RESULT]]
 107 // CHECK_CONSTRAINED:       argpush.block:
 108 // CHECK_CONSTRAINED-NEXT:    [[TMP18:%.*]] = shl i32 [[TMP15]], 2
 109 // CHECK_CONSTRAINED-NEXT:    [[TMP19:%.*]] = or i32 [[TMP18]], 2
 110 // CHECK_CONSTRAINED-NEXT:    store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
 111 // CHECK_CONSTRAINED-NEXT:    [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
 112 // CHECK_CONSTRAINED-NEXT:    store i64 1107004088646384690, ptr addrspace(1) [[TMP20]], align 8
 113 // CHECK_CONSTRAINED-NEXT:    [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8
 114 // CHECK_CONSTRAINED-NEXT:    store i64 8, ptr addrspace(1) [[TMP21]], align 8
 115 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i32 8
 116 // CHECK_CONSTRAINED-NEXT:    store double 3.141590e+00, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
 117 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
 118 // CHECK_CONSTRAINED-NEXT:    store i64 8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8
 119 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8
 120 // CHECK_CONSTRAINED-NEXT:    store i64 4, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8
 121 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8
 122 // CHECK_CONSTRAINED-NEXT:    call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[PRINTBUFFNEXTPTR3]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false)
 123 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i64 [[TMP13]]
 124 // CHECK_CONSTRAINED-NEXT:    store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8
 125 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8
 126 // CHECK_CONSTRAINED-NEXT:    br label [[END_BLOCK]]
 127 //
 128 __device__ int foo1() {
 129   const char *s = "hello world";
 130   return printf("%.*f %*.*s %p\n", 8, 3.14159, 8, 4, s, s);
 131 }
 132
 133 __device__ char *dstr;
 134 __device__ const
 135 // CHECK-LABEL: define dso_local noundef i32 @_Z4foo2v
 136 // CHECK-NEXT:  entry:
 137 // CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
 138 // CHECK-NEXT:    [[LCVAL:%.*]] = alloca i32, align 4, addrspace(5)
 139 // CHECK-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
 140 // CHECK-NEXT:    [[LCVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LCVAL]] to ptr
 141 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8
 142 // CHECK-NEXT:    [[TMP1:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8
 143 // CHECK-NEXT:    [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null
 144 // CHECK-NEXT:    br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
 145 // CHECK:       strlen.while:
 146 // CHECK-NEXT:    [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]
 147 // CHECK-NEXT:    [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1
 148 // CHECK-NEXT:    [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1
 149 // CHECK-NEXT:    [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0
 150 // CHECK-NEXT:    br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
 151 // CHECK:       strlen.while.done:
 152 // CHECK-NEXT:    [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64
 153 // CHECK-NEXT:    [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64
 154 // CHECK-NEXT:    [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]]
 155 // CHECK-NEXT:    [[TMP10:%.*]] = add i64 [[TMP9]], 1
 156 // CHECK-NEXT:    br label [[STRLEN_JOIN]]
 157 // CHECK:       strlen.join:
 158 // CHECK-NEXT:    [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
 159 // CHECK-NEXT:    [[TMP12:%.*]] = add i64 [[TMP11]], 7
 160 // CHECK-NEXT:    [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288
 161 // CHECK-NEXT:    [[TMP14:%.*]] = add i64 [[TMP13]], 36
 162 // CHECK-NEXT:    [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32
 163 // CHECK-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])
 164 // CHECK-NEXT:    [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
 165 // CHECK-NEXT:    br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
 166 // CHECK:       end.block:
 167 // CHECK-NEXT:    [[TMP17:%.*]] = xor i1 [[TMP16]], true
 168 // CHECK-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32
 169 // CHECK-NEXT:    ret i32 [[PRINTF_RESULT]]
 170 // CHECK:       argpush.block:
 171 // CHECK-NEXT:    [[TMP18:%.*]] = shl i32 [[TMP15]], 2
 172 // CHECK-NEXT:    [[TMP19:%.*]] = or i32 [[TMP18]], 2
 173 // CHECK-NEXT:    store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
 174 // CHECK-NEXT:    [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
 175 // CHECK-NEXT:    store i64 7257695813269076350, ptr addrspace(1) [[TMP20]], align 8
 176 // CHECK-NEXT:    [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8
 177 // CHECK-NEXT:    call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP21]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false)
 178 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i64 [[TMP13]]
 179 // CHECK-NEXT:    store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
 180 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
 181 // CHECK-NEXT:    store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo2vE5shval to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8
 182 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8
 183 // CHECK-NEXT:    store ptr [[LCVAL_ASCAST]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8
 184 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8
 185 // CHECK-NEXT:    br label [[END_BLOCK]]
 186 //
 187 // CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo2v
 188 // CHECK_CONSTRAINED-NEXT:  entry:
 189 // CHECK_CONSTRAINED-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
 190 // CHECK_CONSTRAINED-NEXT:    [[LCVAL:%.*]] = alloca i32, align 4, addrspace(5)
 191 // CHECK_CONSTRAINED-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
 192 // CHECK_CONSTRAINED-NEXT:    [[LCVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LCVAL]] to ptr
 193 // CHECK_CONSTRAINED-NEXT:    [[TMP0:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8
 194 // CHECK_CONSTRAINED-NEXT:    [[TMP1:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8
 195 // CHECK_CONSTRAINED-NEXT:    [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null
 196 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
 197 // CHECK_CONSTRAINED:       strlen.while:
 198 // CHECK_CONSTRAINED-NEXT:    [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]
 199 // CHECK_CONSTRAINED-NEXT:    [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1
 200 // CHECK_CONSTRAINED-NEXT:    [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1
 201 // CHECK_CONSTRAINED-NEXT:    [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0
 202 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
 203 // CHECK_CONSTRAINED:       strlen.while.done:
 204 // CHECK_CONSTRAINED-NEXT:    [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64
 205 // CHECK_CONSTRAINED-NEXT:    [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64
 206 // CHECK_CONSTRAINED-NEXT:    [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]]
 207 // CHECK_CONSTRAINED-NEXT:    [[TMP10:%.*]] = add i64 [[TMP9]], 1
 208 // CHECK_CONSTRAINED-NEXT:    br label [[STRLEN_JOIN]]
 209 // CHECK_CONSTRAINED:       strlen.join:
 210 // CHECK_CONSTRAINED-NEXT:    [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
 211 // CHECK_CONSTRAINED-NEXT:    [[TMP12:%.*]] = add i64 [[TMP11]], 7
 212 // CHECK_CONSTRAINED-NEXT:    [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288
 213 // CHECK_CONSTRAINED-NEXT:    [[TMP14:%.*]] = add i64 [[TMP13]], 36
 214 // CHECK_CONSTRAINED-NEXT:    [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32
 215 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])
 216 // CHECK_CONSTRAINED-NEXT:    [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
 217 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
 218 // CHECK_CONSTRAINED:       end.block:
 219 // CHECK_CONSTRAINED-NEXT:    [[TMP17:%.*]] = xor i1 [[TMP16]], true
 220 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32
 221 // CHECK_CONSTRAINED-NEXT:    ret i32 [[PRINTF_RESULT]]
 222 // CHECK_CONSTRAINED:       argpush.block:
 223 // CHECK_CONSTRAINED-NEXT:    [[TMP18:%.*]] = shl i32 [[TMP15]], 2
 224 // CHECK_CONSTRAINED-NEXT:    [[TMP19:%.*]] = or i32 [[TMP18]], 2
 225 // CHECK_CONSTRAINED-NEXT:    store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
 226 // CHECK_CONSTRAINED-NEXT:    [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
 227 // CHECK_CONSTRAINED-NEXT:    store i64 7257695813269076350, ptr addrspace(1) [[TMP20]], align 8
 228 // CHECK_CONSTRAINED-NEXT:    [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8
 229 // CHECK_CONSTRAINED-NEXT:    call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP21]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false)
 230 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i64 [[TMP13]]
 231 // CHECK_CONSTRAINED-NEXT:    store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
 232 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
 233 // CHECK_CONSTRAINED-NEXT:    store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo2vE5shval to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8
 234 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8
 235 // CHECK_CONSTRAINED-NEXT:    store ptr [[LCVAL_ASCAST]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8
 236 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8
 237 // CHECK_CONSTRAINED-NEXT:    br label [[END_BLOCK]]
 238 //
 239 __device__ int foo2() {
 240   __shared__ int shval;
 241   int lcval;
 242   return printf("%s %p %p %p\n", dstr, dstr, &shval, &lcval);
 243 }
 244
 245 __device__ unsigned short g = 30;
 246 __device__ unsigned long n = 30;
 247
 248 __device__ float f1 = 3.14f;
 249 __device__ double f2 = 2.71828;
 250 __device__ _Float16 f3 = 2.71;
 251 __device__ __bf16 f4 = 3.142;
 252 __device__ _BitInt(55) Int55 = 31;
 253 __device__ _BitInt(44) Int44 = 312;
 254 __device__ _BitInt(128) Int128 = 45637;
 255
 256 // CHECK-LABEL: define dso_local noundef i32 @_Z4foo3v
 257 // CHECK-NEXT:  entry:
 258 // CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
 259 // CHECK-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
 260 // CHECK-NEXT:    store i32 25, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4
 261 // CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4
 262 // CHECK-NEXT:    [[TMP1:%.*]] = load i16, ptr addrspacecast (ptr addrspace(1) @g to ptr), align 2
 263 // CHECK-NEXT:    [[CONV:%.*]] = zext i16 [[TMP1]] to i32
 264 // CHECK-NEXT:    [[TMP2:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @n to ptr), align 8
 265 // CHECK-NEXT:    [[TMP3:%.*]] = load float, ptr addrspacecast (ptr addrspace(1) @f1 to ptr), align 4
 266 // CHECK-NEXT:    [[CONV1:%.*]] = fpext float [[TMP3]] to double
 267 // CHECK-NEXT:    [[TMP4:%.*]] = load double, ptr addrspacecast (ptr addrspace(1) @f2 to ptr), align 8
 268 // CHECK-NEXT:    [[TMP5:%.*]] = load half, ptr addrspacecast (ptr addrspace(1) @f3 to ptr), align 2
 269 // CHECK-NEXT:    [[TMP6:%.*]] = load bfloat, ptr addrspacecast (ptr addrspace(1) @f4 to ptr), align 2
 270 // CHECK-NEXT:    [[TMP7:%.*]] = load i55, ptr addrspacecast (ptr addrspace(1) @Int55 to ptr), align 8
 271 // CHECK-NEXT:    [[TMP8:%.*]] = load i44, ptr addrspacecast (ptr addrspace(1) @Int44 to ptr), align 8
 272 // CHECK-NEXT:    [[TMP9:%.*]] = load i128, ptr addrspacecast (ptr addrspace(1) @Int128 to ptr), align 8
 273 // CHECK-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 108)
 274 // CHECK-NEXT:    [[TMP10:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
 275 // CHECK-NEXT:    br i1 [[TMP10]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
 276 // CHECK:       end.block:
 277 // CHECK-NEXT:    [[TMP11:%.*]] = xor i1 [[TMP10]], true
 278 // CHECK-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP11]] to i32
 279 // CHECK-NEXT:    ret i32 [[PRINTF_RESULT]]
 280 // CHECK:       argpush.block:
 281 // CHECK-NEXT:    store i32 434, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
 282 // CHECK-NEXT:    [[TMP12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
 283 // CHECK-NEXT:    store i64 7271852820361268873, ptr addrspace(1) [[TMP12]], align 8
 284 // CHECK-NEXT:    [[TMP13:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP12]], i32 8
 285 // CHECK-NEXT:    [[TMP14:%.*]] = zext i32 [[TMP0]] to i64
 286 // CHECK-NEXT:    store i64 [[TMP14]], ptr addrspace(1) [[TMP13]], align 8
 287 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP13]], i32 8
 288 // CHECK-NEXT:    store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
 289 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
 290 // CHECK-NEXT:    [[TMP15:%.*]] = zext i32 [[CONV]] to i64
 291 // CHECK-NEXT:    store i64 [[TMP15]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8
 292 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8
 293 // CHECK-NEXT:    store i64 [[TMP2]], ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], align 8
 294 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i32 8
 295 // CHECK-NEXT:    store double [[CONV1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8
 296 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8
 297 // CHECK-NEXT:    store double [[TMP4]], ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], align 8
 298 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR6:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], i32 8
 299 // CHECK-NEXT:    [[TMP16:%.*]] = fpext half [[TMP5]] to double
 300 // CHECK-NEXT:    store double [[TMP16]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 8
 301 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR7:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], i32 8
 302 // CHECK-NEXT:    [[TMP17:%.*]] = fpext bfloat [[TMP6]] to double
 303 // CHECK-NEXT:    store double [[TMP17]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 8
 304 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR8:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], i32 8
 305 // CHECK-NEXT:    [[TMP18:%.*]] = zext i55 [[TMP7]] to i64
 306 // CHECK-NEXT:    store i64 [[TMP18]], ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], align 8
 307 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR9:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], i32 8
 308 // CHECK-NEXT:    [[TMP19:%.*]] = zext i44 [[TMP8]] to i64
 309 // CHECK-NEXT:    store i64 [[TMP19]], ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], align 8
 310 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR10:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], i32 8
 311 // CHECK-NEXT:    store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], align 8
 312 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR11:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], i32 16
 313 // CHECK-NEXT:    br label [[END_BLOCK]]
 314 //
 315 // CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo3v
 316 // CHECK_CONSTRAINED-NEXT:  entry:
 317 // CHECK_CONSTRAINED-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
 318 // CHECK_CONSTRAINED-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
 319 // CHECK_CONSTRAINED-NEXT:    store i32 25, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4
 320 // CHECK_CONSTRAINED-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4
 321 // CHECK_CONSTRAINED-NEXT:    [[TMP1:%.*]] = load i16, ptr addrspacecast (ptr addrspace(1) @g to ptr), align 2
 322 // CHECK_CONSTRAINED-NEXT:    [[CONV:%.*]] = zext i16 [[TMP1]] to i32
 323 // CHECK_CONSTRAINED-NEXT:    [[TMP2:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @n to ptr), align 8
 324 // CHECK_CONSTRAINED-NEXT:    [[TMP3:%.*]] = load float, ptr addrspacecast (ptr addrspace(1) @f1 to ptr), align 4
 325 // CHECK_CONSTRAINED-NEXT:    [[CONV1:%.*]] = fpext float [[TMP3]] to double
 326 // CHECK_CONSTRAINED-NEXT:    [[TMP4:%.*]] = load double, ptr addrspacecast (ptr addrspace(1) @f2 to ptr), align 8
 327 // CHECK_CONSTRAINED-NEXT:    [[TMP5:%.*]] = load half, ptr addrspacecast (ptr addrspace(1) @f3 to ptr), align 2
 328 // CHECK_CONSTRAINED-NEXT:    [[TMP6:%.*]] = load bfloat, ptr addrspacecast (ptr addrspace(1) @f4 to ptr), align 2
 329 // CHECK_CONSTRAINED-NEXT:    [[TMP7:%.*]] = load i55, ptr addrspacecast (ptr addrspace(1) @Int55 to ptr), align 8
 330 // CHECK_CONSTRAINED-NEXT:    [[TMP8:%.*]] = load i44, ptr addrspacecast (ptr addrspace(1) @Int44 to ptr), align 8
 331 // CHECK_CONSTRAINED-NEXT:    [[TMP9:%.*]] = load i128, ptr addrspacecast (ptr addrspace(1) @Int128 to ptr), align 8
 332 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 108)
 333 // CHECK_CONSTRAINED-NEXT:    [[TMP10:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
 334 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP10]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
 335 // CHECK_CONSTRAINED:       end.block:
 336 // CHECK_CONSTRAINED-NEXT:    [[TMP11:%.*]] = xor i1 [[TMP10]], true
 337 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP11]] to i32
 338 // CHECK_CONSTRAINED-NEXT:    ret i32 [[PRINTF_RESULT]]
 339 // CHECK_CONSTRAINED:       argpush.block:
 340 // CHECK_CONSTRAINED-NEXT:    store i32 434, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
 341 // CHECK_CONSTRAINED-NEXT:    [[TMP12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
 342 // CHECK_CONSTRAINED-NEXT:    store i64 7271852820361268873, ptr addrspace(1) [[TMP12]], align 8
 343 // CHECK_CONSTRAINED-NEXT:    [[TMP13:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP12]], i32 8
 344 // CHECK_CONSTRAINED-NEXT:    [[TMP14:%.*]] = zext i32 [[TMP0]] to i64
 345 // CHECK_CONSTRAINED-NEXT:    store i64 [[TMP14]], ptr addrspace(1) [[TMP13]], align 8
 346 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP13]], i32 8
 347 // CHECK_CONSTRAINED-NEXT:    store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
 348 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
 349 // CHECK_CONSTRAINED-NEXT:    [[TMP15:%.*]] = zext i32 [[CONV]] to i64
 350 // CHECK_CONSTRAINED-NEXT:    store i64 [[TMP15]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8
 351 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8
 352 // CHECK_CONSTRAINED-NEXT:    store i64 [[TMP2]], ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], align 8
 353 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i32 8
 354 // CHECK_CONSTRAINED-NEXT:    store double [[CONV1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8
 355 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8
 356 // CHECK_CONSTRAINED-NEXT:    store double [[TMP4]], ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], align 8
 357 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR6:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], i32 8
 358 // CHECK_CONSTRAINED-NEXT:    [[TMP16:%.*]] = fpext half [[TMP5]] to double
 359 // CHECK_CONSTRAINED-NEXT:    store double [[TMP16]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 8
 360 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR7:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], i32 8
 361 // CHECK_CONSTRAINED-NEXT:    [[TMP17:%.*]] = fpext bfloat [[TMP6]] to double
 362 // CHECK_CONSTRAINED-NEXT:    store double [[TMP17]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 8
 363 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR8:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], i32 8
 364 // CHECK_CONSTRAINED-NEXT:    [[TMP18:%.*]] = zext i55 [[TMP7]] to i64
 365 // CHECK_CONSTRAINED-NEXT:    store i64 [[TMP18]], ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], align 8
 366 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR9:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], i32 8
 367 // CHECK_CONSTRAINED-NEXT:    [[TMP19:%.*]] = zext i44 [[TMP8]] to i64
 368 // CHECK_CONSTRAINED-NEXT:    store i64 [[TMP19]], ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], align 8
 369 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR10:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], i32 8
 370 // CHECK_CONSTRAINED-NEXT:    store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], align 8
 371 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR11:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], i32 16
 372 // CHECK_CONSTRAINED-NEXT:    br label [[END_BLOCK]]
 373 //
 374 __device__ int foo3() {
 375   __shared__ int s;
 376   s = 25;
 377   return printf("Random values: %d,%p,%hd,%ld,%f,%f,%f,%f,%d,%d,%d\n",s, &s, g, n, f1, f2, f3, f4, Int55, Int44, Int128);
 378 }
 379
 380 //A non trivial case,
 381 // CHECK-LABEL: define dso_local noundef i32 @_Z4foo4v
 382 // CHECK-NEXT:  entry:
 383 // CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
 384 // CHECK-NEXT:    [[S:%.*]] = alloca ptr, align 8, addrspace(5)
 385 // CHECK-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
 386 // CHECK-NEXT:    [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr
 387 // CHECK-NEXT:    store ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr), ptr [[S_ASCAST]], align 8
 388 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
 389 // CHECK-NEXT:    [[TMP1:%.*]] = icmp eq ptr [[TMP0]], null
 390 // CHECK-NEXT:    br i1 [[TMP1]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
 391 // CHECK:       strlen.while:
 392 // CHECK-NEXT:    [[TMP2:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP3:%.*]], [[STRLEN_WHILE]] ]
 393 // CHECK-NEXT:    [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
 394 // CHECK-NEXT:    [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
 395 // CHECK-NEXT:    [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
 396 // CHECK-NEXT:    br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
 397 // CHECK:       strlen.while.done:
 398 // CHECK-NEXT:    [[TMP6:%.*]] = ptrtoint ptr [[TMP0]] to i64
 399 // CHECK-NEXT:    [[TMP7:%.*]] = ptrtoint ptr [[TMP2]] to i64
 400 // CHECK-NEXT:    [[TMP8:%.*]] = sub i64 [[TMP7]], [[TMP6]]
 401 // CHECK-NEXT:    [[TMP9:%.*]] = add i64 [[TMP8]], 1
 402 // CHECK-NEXT:    br label [[STRLEN_JOIN]]
 403 // CHECK:       strlen.join:
 404 // CHECK-NEXT:    [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
 405 // CHECK-NEXT:    [[TMP11:%.*]] = add i64 [[TMP10]], 7
 406 // CHECK-NEXT:    [[TMP12:%.*]] = and i64 [[TMP11]], 4294967288
 407 // CHECK-NEXT:    [[TMP13:%.*]] = add i64 [[TMP12]], 12
 408 // CHECK-NEXT:    [[TMP14:%.*]] = trunc i64 [[TMP13]] to i32
 409 // CHECK-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP14]])
 410 // CHECK-NEXT:    [[TMP15:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
 411 // CHECK-NEXT:    br i1 [[TMP15]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
 412 // CHECK:       end.block:
 413 // CHECK-NEXT:    [[TMP16:%.*]] = xor i1 [[TMP15]], true
 414 // CHECK-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP16]] to i32
 415 // CHECK-NEXT:    ret i32 [[PRINTF_RESULT]]
 416 // CHECK:       argpush.block:
 417 // CHECK-NEXT:    [[TMP17:%.*]] = shl i32 [[TMP14]], 2
 418 // CHECK-NEXT:    store i32 [[TMP17]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
 419 // CHECK-NEXT:    [[TMP18:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
 420 // CHECK-NEXT:    call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP18]], ptr align 1 [[TMP0]], i64 [[TMP10]], i1 false)
 421 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP18]], i64 [[TMP12]]
 422 // CHECK-NEXT:    store i64 10, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
 423 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
 424 // CHECK-NEXT:    br label [[END_BLOCK]]
 425 //
 426 // CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo4v
 427 // CHECK_CONSTRAINED-NEXT:  entry:
 428 // CHECK_CONSTRAINED-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
 429 // CHECK_CONSTRAINED-NEXT:    [[S:%.*]] = alloca ptr, align 8, addrspace(5)
 430 // CHECK_CONSTRAINED-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
 431 // CHECK_CONSTRAINED-NEXT:    [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr
 432 // CHECK_CONSTRAINED-NEXT:    store ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr), ptr [[S_ASCAST]], align 8
 433 // CHECK_CONSTRAINED-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
 434 // CHECK_CONSTRAINED-NEXT:    [[TMP1:%.*]] = icmp eq ptr [[TMP0]], null
 435 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP1]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
 436 // CHECK_CONSTRAINED:       strlen.while:
 437 // CHECK_CONSTRAINED-NEXT:    [[TMP2:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP3:%.*]], [[STRLEN_WHILE]] ]
 438 // CHECK_CONSTRAINED-NEXT:    [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
 439 // CHECK_CONSTRAINED-NEXT:    [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
 440 // CHECK_CONSTRAINED-NEXT:    [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
 441 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
 442 // CHECK_CONSTRAINED:       strlen.while.done:
 443 // CHECK_CONSTRAINED-NEXT:    [[TMP6:%.*]] = ptrtoint ptr [[TMP0]] to i64
 444 // CHECK_CONSTRAINED-NEXT:    [[TMP7:%.*]] = ptrtoint ptr [[TMP2]] to i64
 445 // CHECK_CONSTRAINED-NEXT:    [[TMP8:%.*]] = sub i64 [[TMP7]], [[TMP6]]
 446 // CHECK_CONSTRAINED-NEXT:    [[TMP9:%.*]] = add i64 [[TMP8]], 1
 447 // CHECK_CONSTRAINED-NEXT:    br label [[STRLEN_JOIN]]
 448 // CHECK_CONSTRAINED:       strlen.join:
 449 // CHECK_CONSTRAINED-NEXT:    [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
 450 // CHECK_CONSTRAINED-NEXT:    [[TMP11:%.*]] = add i64 [[TMP10]], 7
 451 // CHECK_CONSTRAINED-NEXT:    [[TMP12:%.*]] = and i64 [[TMP11]], 4294967288
 452 // CHECK_CONSTRAINED-NEXT:    [[TMP13:%.*]] = add i64 [[TMP12]], 12
 453 // CHECK_CONSTRAINED-NEXT:    [[TMP14:%.*]] = trunc i64 [[TMP13]] to i32
 454 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP14]])
 455 // CHECK_CONSTRAINED-NEXT:    [[TMP15:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
 456 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP15]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
 457 // CHECK_CONSTRAINED:       end.block:
 458 // CHECK_CONSTRAINED-NEXT:    [[TMP16:%.*]] = xor i1 [[TMP15]], true
 459 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP16]] to i32
 460 // CHECK_CONSTRAINED-NEXT:    ret i32 [[PRINTF_RESULT]]
 461 // CHECK_CONSTRAINED:       argpush.block:
 462 // CHECK_CONSTRAINED-NEXT:    [[TMP17:%.*]] = shl i32 [[TMP14]], 2
 463 // CHECK_CONSTRAINED-NEXT:    store i32 [[TMP17]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
 464 // CHECK_CONSTRAINED-NEXT:    [[TMP18:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
 465 // CHECK_CONSTRAINED-NEXT:    call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP18]], ptr align 1 [[TMP0]], i64 [[TMP10]], i1 false)
 466 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP18]], i64 [[TMP12]]
 467 // CHECK_CONSTRAINED-NEXT:    store i64 10, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
 468 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
 469 // CHECK_CONSTRAINED-NEXT:    br label [[END_BLOCK]]
 470 //
 471 __device__ int foo4() {
 472   const char* s = "format str%d";
 473   return printf(s, 10);
 474 }