xref: /llvm-project/clang/test/CodeGenHIP/printf_nonhostcall.cpp (revision ecbe4d1e360e25c0634a3a62fbd01e8df5bb0c1b) 
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 contract 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 i64, ptr addrspacecast (ptr addrspace(1) @Int55 to ptr), align 8
271 // CHECK-NEXT:    [[LOADEDV:%.*]] = trunc i64 [[TMP7]] to i55
272 // CHECK-NEXT:    [[TMP8:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @Int44 to ptr), align 8
273 // CHECK-NEXT:    [[LOADEDV2:%.*]] = trunc i64 [[TMP8]] to i44
274 // CHECK-NEXT:    [[TMP9:%.*]] = load i128, ptr addrspacecast (ptr addrspace(1) @Int128 to ptr), align 8
275 // CHECK-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 108)
276 // CHECK-NEXT:    [[TMP10:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
277 // CHECK-NEXT:    br i1 [[TMP10]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
278 // CHECK:       end.block:
279 // CHECK-NEXT:    [[TMP11:%.*]] = xor i1 [[TMP10]], true
280 // CHECK-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP11]] to i32
281 // CHECK-NEXT:    ret i32 [[PRINTF_RESULT]]
282 // CHECK:       argpush.block:
283 // CHECK-NEXT:    store i32 434, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
284 // CHECK-NEXT:    [[TMP12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
285 // CHECK-NEXT:    store i64 7271852820361268873, ptr addrspace(1) [[TMP12]], align 8
286 // CHECK-NEXT:    [[TMP13:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP12]], i32 8
287 // CHECK-NEXT:    [[TMP14:%.*]] = zext i32 [[TMP0]] to i64
288 // CHECK-NEXT:    store i64 [[TMP14]], ptr addrspace(1) [[TMP13]], align 8
289 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP13]], i32 8
290 // CHECK-NEXT:    store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
291 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
292 // CHECK-NEXT:    [[TMP15:%.*]] = zext i32 [[CONV]] to i64
293 // CHECK-NEXT:    store i64 [[TMP15]], ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], align 8
294 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i32 8
295 // CHECK-NEXT:    store i64 [[TMP2]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8
296 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8
297 // CHECK-NEXT:    store double [[CONV1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], align 8
298 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR6:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], i32 8
299 // CHECK-NEXT:    store double [[TMP4]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 8
300 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR7:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], i32 8
301 // CHECK-NEXT:    [[TMP16:%.*]] = fpext half [[TMP5]] to double
302 // CHECK-NEXT:    store double [[TMP16]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 8
303 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR8:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], i32 8
304 // CHECK-NEXT:    [[TMP17:%.*]] = fpext bfloat [[TMP6]] to double
305 // CHECK-NEXT:    store double [[TMP17]], ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], align 8
306 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR9:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], i32 8
307 // CHECK-NEXT:    [[TMP18:%.*]] = zext i55 [[LOADEDV]] to i64
308 // CHECK-NEXT:    store i64 [[TMP18]], ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], align 8
309 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR10:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], i32 8
310 // CHECK-NEXT:    [[TMP19:%.*]] = zext i44 [[LOADEDV2]] to i64
311 // CHECK-NEXT:    store i64 [[TMP19]], ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], align 8
312 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR11:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], i32 8
313 // CHECK-NEXT:    store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], align 8
314 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], i32 16
315 // CHECK-NEXT:    br label [[END_BLOCK]]
316 //
317 // CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo3v
318 // CHECK_CONSTRAINED-NEXT:  entry:
319 // CHECK_CONSTRAINED-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
320 // CHECK_CONSTRAINED-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
321 // CHECK_CONSTRAINED-NEXT:    store i32 25, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4
322 // CHECK_CONSTRAINED-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4
323 // CHECK_CONSTRAINED-NEXT:    [[TMP1:%.*]] = load i16, ptr addrspacecast (ptr addrspace(1) @g to ptr), align 2
324 // CHECK_CONSTRAINED-NEXT:    [[CONV:%.*]] = zext i16 [[TMP1]] to i32
325 // CHECK_CONSTRAINED-NEXT:    [[TMP2:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @n to ptr), align 8
326 // CHECK_CONSTRAINED-NEXT:    [[TMP3:%.*]] = load float, ptr addrspacecast (ptr addrspace(1) @f1 to ptr), align 4
327 // CHECK_CONSTRAINED-NEXT:    [[CONV1:%.*]] = fpext contract float [[TMP3]] to double
328 // CHECK_CONSTRAINED-NEXT:    [[TMP4:%.*]] = load double, ptr addrspacecast (ptr addrspace(1) @f2 to ptr), align 8
329 // CHECK_CONSTRAINED-NEXT:    [[TMP5:%.*]] = load half, ptr addrspacecast (ptr addrspace(1) @f3 to ptr), align 2
330 // CHECK_CONSTRAINED-NEXT:    [[TMP6:%.*]] = load bfloat, ptr addrspacecast (ptr addrspace(1) @f4 to ptr), align 2
331 // CHECK_CONSTRAINED-NEXT:    [[TMP7:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @Int55 to ptr), align 8
332 // CHECK_CONSTRAINED-NEXT:    [[LOADEDV:%.*]] = trunc i64 [[TMP7]] to i55
333 // CHECK_CONSTRAINED-NEXT:    [[TMP8:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @Int44 to ptr), align 8
334 // CHECK_CONSTRAINED-NEXT:    [[LOADEDV2:%.*]] = trunc i64 [[TMP8]] to i44
335 // CHECK_CONSTRAINED-NEXT:    [[TMP9:%.*]] = load i128, ptr addrspacecast (ptr addrspace(1) @Int128 to ptr), align 8
336 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 108)
337 // CHECK_CONSTRAINED-NEXT:    [[TMP10:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
338 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP10]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
339 // CHECK_CONSTRAINED:       end.block:
340 // CHECK_CONSTRAINED-NEXT:    [[TMP11:%.*]] = xor i1 [[TMP10]], true
341 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP11]] to i32
342 // CHECK_CONSTRAINED-NEXT:    ret i32 [[PRINTF_RESULT]]
343 // CHECK_CONSTRAINED:       argpush.block:
344 // CHECK_CONSTRAINED-NEXT:    store i32 434, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
345 // CHECK_CONSTRAINED-NEXT:    [[TMP12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
346 // CHECK_CONSTRAINED-NEXT:    store i64 7271852820361268873, ptr addrspace(1) [[TMP12]], align 8
347 // CHECK_CONSTRAINED-NEXT:    [[TMP13:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP12]], i32 8
348 // CHECK_CONSTRAINED-NEXT:    [[TMP14:%.*]] = zext i32 [[TMP0]] to i64
349 // CHECK_CONSTRAINED-NEXT:    store i64 [[TMP14]], ptr addrspace(1) [[TMP13]], align 8
350 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP13]], i32 8
351 // CHECK_CONSTRAINED-NEXT:    store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
352 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
353 // CHECK_CONSTRAINED-NEXT:    [[TMP15:%.*]] = zext i32 [[CONV]] to i64
354 // CHECK_CONSTRAINED-NEXT:    store i64 [[TMP15]], ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], align 8
355 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i32 8
356 // CHECK_CONSTRAINED-NEXT:    store i64 [[TMP2]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8
357 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8
358 // CHECK_CONSTRAINED-NEXT:    store double [[CONV1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], align 8
359 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR6:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], i32 8
360 // CHECK_CONSTRAINED-NEXT:    store double [[TMP4]], ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], align 8
361 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR7:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR6]], i32 8
362 // CHECK_CONSTRAINED-NEXT:    [[TMP16:%.*]] = fpext half [[TMP5]] to double
363 // CHECK_CONSTRAINED-NEXT:    store double [[TMP16]], ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], align 8
364 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR8:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR7]], i32 8
365 // CHECK_CONSTRAINED-NEXT:    [[TMP17:%.*]] = fpext bfloat [[TMP6]] to double
366 // CHECK_CONSTRAINED-NEXT:    store double [[TMP17]], ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], align 8
367 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR9:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR8]], i32 8
368 // CHECK_CONSTRAINED-NEXT:    [[TMP18:%.*]] = zext i55 [[LOADEDV]] to i64
369 // CHECK_CONSTRAINED-NEXT:    store i64 [[TMP18]], ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], align 8
370 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR10:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR9]], i32 8
371 // CHECK_CONSTRAINED-NEXT:    [[TMP19:%.*]] = zext i44 [[LOADEDV2]] to i64
372 // CHECK_CONSTRAINED-NEXT:    store i64 [[TMP19]], ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], align 8
373 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR11:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR10]], i32 8
374 // CHECK_CONSTRAINED-NEXT:    store i128 [[TMP9]], ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], align 8
375 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR12:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR11]], i32 16
376 // CHECK_CONSTRAINED-NEXT:    br label [[END_BLOCK]]
377 //
378 __device__ int foo3() {
379   __shared__ int s;
380   s = 25;
381   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);
382 }
383 
384 //A non trivial case,
385 // CHECK-LABEL: define dso_local noundef i32 @_Z4foo4v
386 // CHECK-NEXT:  entry:
387 // CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
388 // CHECK-NEXT:    [[S:%.*]] = alloca ptr, align 8, addrspace(5)
389 // CHECK-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
390 // CHECK-NEXT:    [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr
391 // CHECK-NEXT:    store ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr), ptr [[S_ASCAST]], align 8
392 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
393 // CHECK-NEXT:    [[TMP1:%.*]] = icmp eq ptr [[TMP0]], null
394 // CHECK-NEXT:    br i1 [[TMP1]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
395 // CHECK:       strlen.while:
396 // CHECK-NEXT:    [[TMP2:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP3:%.*]], [[STRLEN_WHILE]] ]
397 // CHECK-NEXT:    [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
398 // CHECK-NEXT:    [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
399 // CHECK-NEXT:    [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
400 // CHECK-NEXT:    br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
401 // CHECK:       strlen.while.done:
402 // CHECK-NEXT:    [[TMP6:%.*]] = ptrtoint ptr [[TMP0]] to i64
403 // CHECK-NEXT:    [[TMP7:%.*]] = ptrtoint ptr [[TMP2]] to i64
404 // CHECK-NEXT:    [[TMP8:%.*]] = sub i64 [[TMP7]], [[TMP6]]
405 // CHECK-NEXT:    [[TMP9:%.*]] = add i64 [[TMP8]], 1
406 // CHECK-NEXT:    br label [[STRLEN_JOIN]]
407 // CHECK:       strlen.join:
408 // CHECK-NEXT:    [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
409 // CHECK-NEXT:    [[TMP11:%.*]] = add i64 [[TMP10]], 7
410 // CHECK-NEXT:    [[TMP12:%.*]] = and i64 [[TMP11]], 4294967288
411 // CHECK-NEXT:    [[TMP13:%.*]] = add i64 [[TMP12]], 12
412 // CHECK-NEXT:    [[TMP14:%.*]] = trunc i64 [[TMP13]] to i32
413 // CHECK-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP14]])
414 // CHECK-NEXT:    [[TMP15:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
415 // CHECK-NEXT:    br i1 [[TMP15]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
416 // CHECK:       end.block:
417 // CHECK-NEXT:    [[TMP16:%.*]] = xor i1 [[TMP15]], true
418 // CHECK-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP16]] to i32
419 // CHECK-NEXT:    ret i32 [[PRINTF_RESULT]]
420 // CHECK:       argpush.block:
421 // CHECK-NEXT:    [[TMP17:%.*]] = shl i32 [[TMP14]], 2
422 // CHECK-NEXT:    store i32 [[TMP17]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
423 // CHECK-NEXT:    [[TMP18:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
424 // CHECK-NEXT:    call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP18]], ptr align 1 [[TMP0]], i64 [[TMP10]], i1 false)
425 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP18]], i64 [[TMP12]]
426 // CHECK-NEXT:    store i64 10, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
427 // CHECK-NEXT:    [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
428 // CHECK-NEXT:    br label [[END_BLOCK]]
429 //
430 // CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo4v
431 // CHECK_CONSTRAINED-NEXT:  entry:
432 // CHECK_CONSTRAINED-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
433 // CHECK_CONSTRAINED-NEXT:    [[S:%.*]] = alloca ptr, align 8, addrspace(5)
434 // CHECK_CONSTRAINED-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
435 // CHECK_CONSTRAINED-NEXT:    [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr
436 // CHECK_CONSTRAINED-NEXT:    store ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr), ptr [[S_ASCAST]], align 8
437 // CHECK_CONSTRAINED-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
438 // CHECK_CONSTRAINED-NEXT:    [[TMP1:%.*]] = icmp eq ptr [[TMP0]], null
439 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP1]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
440 // CHECK_CONSTRAINED:       strlen.while:
441 // CHECK_CONSTRAINED-NEXT:    [[TMP2:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP3:%.*]], [[STRLEN_WHILE]] ]
442 // CHECK_CONSTRAINED-NEXT:    [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
443 // CHECK_CONSTRAINED-NEXT:    [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
444 // CHECK_CONSTRAINED-NEXT:    [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
445 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
446 // CHECK_CONSTRAINED:       strlen.while.done:
447 // CHECK_CONSTRAINED-NEXT:    [[TMP6:%.*]] = ptrtoint ptr [[TMP0]] to i64
448 // CHECK_CONSTRAINED-NEXT:    [[TMP7:%.*]] = ptrtoint ptr [[TMP2]] to i64
449 // CHECK_CONSTRAINED-NEXT:    [[TMP8:%.*]] = sub i64 [[TMP7]], [[TMP6]]
450 // CHECK_CONSTRAINED-NEXT:    [[TMP9:%.*]] = add i64 [[TMP8]], 1
451 // CHECK_CONSTRAINED-NEXT:    br label [[STRLEN_JOIN]]
452 // CHECK_CONSTRAINED:       strlen.join:
453 // CHECK_CONSTRAINED-NEXT:    [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
454 // CHECK_CONSTRAINED-NEXT:    [[TMP11:%.*]] = add i64 [[TMP10]], 7
455 // CHECK_CONSTRAINED-NEXT:    [[TMP12:%.*]] = and i64 [[TMP11]], 4294967288
456 // CHECK_CONSTRAINED-NEXT:    [[TMP13:%.*]] = add i64 [[TMP12]], 12
457 // CHECK_CONSTRAINED-NEXT:    [[TMP14:%.*]] = trunc i64 [[TMP13]] to i32
458 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP14]])
459 // CHECK_CONSTRAINED-NEXT:    [[TMP15:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
460 // CHECK_CONSTRAINED-NEXT:    br i1 [[TMP15]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
461 // CHECK_CONSTRAINED:       end.block:
462 // CHECK_CONSTRAINED-NEXT:    [[TMP16:%.*]] = xor i1 [[TMP15]], true
463 // CHECK_CONSTRAINED-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP16]] to i32
464 // CHECK_CONSTRAINED-NEXT:    ret i32 [[PRINTF_RESULT]]
465 // CHECK_CONSTRAINED:       argpush.block:
466 // CHECK_CONSTRAINED-NEXT:    [[TMP17:%.*]] = shl i32 [[TMP14]], 2
467 // CHECK_CONSTRAINED-NEXT:    store i32 [[TMP17]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
468 // CHECK_CONSTRAINED-NEXT:    [[TMP18:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
469 // CHECK_CONSTRAINED-NEXT:    call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP18]], ptr align 1 [[TMP0]], i64 [[TMP10]], i1 false)
470 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP18]], i64 [[TMP12]]
471 // CHECK_CONSTRAINED-NEXT:    store i64 10, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
472 // CHECK_CONSTRAINED-NEXT:    [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
473 // CHECK_CONSTRAINED-NEXT:    br label [[END_BLOCK]]
474 //
475 __device__ int foo4() {
476   const char* s = "format str%d";
477   return printf(s, 10);
478 }
479