| blob | 69efc47008e6aade644fb8684c536831257042c5 |
1 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK %s
2 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx802 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK %s
3 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK %s
4 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
5 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx802 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
6 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
8 %struct.A = type { i8, float }
9 %opencl.image1d_t = type opaque
10 %opencl.image2d_t = type opaque
11 %opencl.image3d_t = type opaque
12 %opencl.queue_t = type opaque
13 %opencl.pipe_t = type opaque
14 %struct.B = type { ptr addrspace(1) }
15 %opencl.clk_event_t = type opaque
17 @__test_block_invoke_kernel_runtime_handle = external addrspace(1) externally_initialized constant ptr addrspace(1)
19 ; CHECK: ---
20 ; CHECK-NEXT: amdhsa.kernels:
21 ; CHECK-NEXT: - .args:
22 ; CHECK-NEXT: - .name: a
23 ; CHECK-NEXT: .offset: 0
24 ; CHECK-NEXT: .size: 1
25 ; CHECK-NEXT: .type_name: char
26 ; CHECK-NEXT: .value_kind: by_value
27 ; CHECK-NEXT: - .offset: 8
28 ; CHECK-NEXT: .size: 8
29 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
30 ; CHECK-NEXT: - .offset: 16
31 ; CHECK-NEXT: .size: 8
32 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
33 ; CHECK-NEXT: - .offset: 24
34 ; CHECK-NEXT: .size: 8
35 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
36 ; CHECK-NEXT: - .offset: 32
37 ; CHECK-NEXT: .size: 8
38 ; CHECK-NOT: .value_kind: hidden_default_queue
39 ; CHECK-NOT: .value_kind: hidden_completion_action
40 ; CHECK-NOT: .value_kind: hidden_hostcall_buffer
41 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
42 ; CHECK: .value_kind: hidden_multigrid_sync_arg
43 ; CHECK: .language: OpenCL C
44 ; CHECK-NEXT: .language_version:
45 ; CHECK-NEXT: - 2
46 ; CHECK-NEXT: - 0
47 ; CHECK: .name: test_char
48 ; CHECK: .symbol: test_char.kd
49 define amdgpu_kernel void @test_char(i8 %a) #0
50 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
51 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
52 ret void
53 }
55 ; CHECK: - .args:
56 ; CHECK-NEXT: - .name: a
57 ; CHECK-NEXT: .offset: 0
58 ; CHECK-NEXT: .size: 1
59 ; CHECK-NEXT: .type_name: char
60 ; CHECK-NEXT: .value_kind: by_value
61 ; CHECK-NEXT: - .offset: 8
62 ; CHECK-NEXT: .size: 8
63 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
64 ; CHECK-NEXT: - .offset: 16
65 ; CHECK-NEXT: .size: 8
66 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
67 ; CHECK-NEXT: - .offset: 24
68 ; CHECK-NEXT: .size: 8
69 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
70 ; CHECK-NEXT: - .offset: 32
71 ; CHECK-NEXT: .size: 8
72 ; CHECK-NOT: .value_kind: hidden_default_queue
73 ; CHECK-NOT: .value_kind: hidden_completion_action
74 ; CHECK-NOT: .value_kind: hidden_hostcall_buffer
75 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
76 ; CHECK: .value_kind: hidden_multigrid_sync_arg
77 ; CHECK: .language: OpenCL C
78 ; CHECK-NEXT: .language_version:
79 ; CHECK-NEXT: - 2
80 ; CHECK-NEXT: - 0
81 ; CHECK: .name: test_char_byref_constant
82 ; CHECK: .symbol: test_char_byref_constant.kd
83 define amdgpu_kernel void @test_char_byref_constant(ptr addrspace(4) byref(i8) %a) #0
84 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
85 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
86 ret void
87 }
89 ; CHECK: - .args:
90 ; CHECK-NEXT: - .offset: 0
91 ; CHECK-NEXT: .size: 1
92 ; CHECK-NEXT: .type_name: char
93 ; CHECK-NEXT: .value_kind: by_value
94 ; CHECK-NEXT: - .name: a
95 ; CHECK-NEXT: .offset: 512
96 ; CHECK-NEXT: .size: 1
97 ; CHECK-NEXT: .type_name: char
98 ; CHECK-NEXT: .value_kind: by_value
99 ; CHECK-NEXT: - .offset: 520
100 ; CHECK-NEXT: .size: 8
101 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
102 ; CHECK-NEXT: - .offset: 528
103 ; CHECK-NEXT: .size: 8
104 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
105 ; CHECK-NEXT: - .offset: 536
106 ; CHECK-NEXT: .size: 8
107 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
108 ; CHECK-NEXT: - .offset: 544
109 ; CHECK-NEXT: .size: 8
110 ; CHECK-NOT: .value_kind: hidden_default_queue
111 ; CHECK-NOT: .value_kind: hidden_completion_action
112 ; CHECK-NOT: .value_kind: hidden_hostcall_buffer
113 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
114 ; CHECK: .value_kind: hidden_multigrid_sync_arg
115 ; CHECK: .language: OpenCL C
116 ; CHECK-NEXT: .language_version:
117 ; CHECK-NEXT: - 2
118 ; CHECK-NEXT: - 0
119 ; CHECK: .name: test_char_byref_constant_align512
120 ; CHECK: .symbol: test_char_byref_constant_align512.kd
121 define amdgpu_kernel void @test_char_byref_constant_align512(i8, ptr addrspace(4) byref(i8) align(512) %a) #0
122 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !111
123 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
124 ret void
125 }
127 ; CHECK: - .args:
128 ; CHECK-NEXT: - .name: a
129 ; CHECK-NEXT: .offset: 0
130 ; CHECK-NEXT: .size: 4
131 ; CHECK-NEXT: .type_name: ushort2
132 ; CHECK-NEXT: .value_kind: by_value
133 ; CHECK-NEXT: - .offset: 8
134 ; CHECK-NEXT: .size: 8
135 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
136 ; CHECK-NEXT: - .offset: 16
137 ; CHECK-NEXT: .size: 8
138 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
139 ; CHECK-NEXT: - .offset: 24
140 ; CHECK-NEXT: .size: 8
141 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
142 ; CHECK-NEXT: - .offset: 32
143 ; CHECK-NEXT: .size: 8
144 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
145 ; CHECK-NEXT: - .offset: 40
146 ; CHECK-NEXT: .size: 8
147 ; CHECK-NEXT: .value_kind: hidden_none
148 ; CHECK-NEXT: - .offset: 48
149 ; CHECK-NEXT: .size: 8
150 ; CHECK-NEXT: .value_kind: hidden_none
151 ; CHECK-NEXT: - .offset: 56
152 ; CHECK-NEXT: .size: 8
153 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
154 ; CHECK: .language: OpenCL C
155 ; CHECK-NEXT: .language_version:
156 ; CHECK-NEXT: - 2
157 ; CHECK-NEXT: - 0
158 ; CHECK: .name: test_ushort2
159 ; CHECK: .symbol: test_ushort2.kd
160 define amdgpu_kernel void @test_ushort2(<2 x i16> %a) #0
161 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !10
162 !kernel_arg_base_type !10 !kernel_arg_type_qual !4 {
163 ret void
164 }
166 ; CHECK: - .args:
167 ; CHECK-NEXT: - .name: a
168 ; CHECK-NEXT: .offset: 0
169 ; CHECK-NEXT: .size: 16
170 ; CHECK-NEXT: .type_name: int3
171 ; CHECK-NEXT: .value_kind: by_value
172 ; CHECK-NEXT: - .offset: 16
173 ; CHECK-NEXT: .size: 8
174 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
175 ; CHECK-NEXT: - .offset: 24
176 ; CHECK-NEXT: .size: 8
177 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
178 ; CHECK-NEXT: - .offset: 32
179 ; CHECK-NEXT: .size: 8
180 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
181 ; CHECK-NEXT: - .offset: 40
182 ; CHECK-NEXT: .size: 8
183 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
184 ; CHECK-NEXT: - .offset: 48
185 ; CHECK-NEXT: .size: 8
186 ; CHECK-NEXT: .value_kind: hidden_none
187 ; CHECK-NEXT: - .offset: 56
188 ; CHECK-NEXT: .size: 8
189 ; CHECK-NEXT: .value_kind: hidden_none
190 ; CHECK-NEXT: - .offset: 64
191 ; CHECK-NEXT: .size: 8
192 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
193 ; CHECK: .language: OpenCL C
194 ; CHECK-NEXT: .language_version:
195 ; CHECK-NEXT: - 2
196 ; CHECK-NEXT: - 0
197 ; CHECK: .name: test_int3
198 ; CHECK: .symbol: test_int3.kd
199 define amdgpu_kernel void @test_int3(<3 x i32> %a) #0
200 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !11
201 !kernel_arg_base_type !11 !kernel_arg_type_qual !4 {
202 ret void
203 }
205 ; CHECK: - .args:
206 ; CHECK-NEXT: - .name: a
207 ; CHECK-NEXT: .offset: 0
208 ; CHECK-NEXT: .size: 32
209 ; CHECK-NEXT: .type_name: ulong4
210 ; CHECK-NEXT: .value_kind: by_value
211 ; CHECK-NEXT: - .offset: 32
212 ; CHECK-NEXT: .size: 8
213 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
214 ; CHECK-NEXT: - .offset: 40
215 ; CHECK-NEXT: .size: 8
216 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
217 ; CHECK-NEXT: - .offset: 48
218 ; CHECK-NEXT: .size: 8
219 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
220 ; CHECK-NEXT: - .offset: 56
221 ; CHECK-NEXT: .size: 8
222 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
223 ; CHECK-NEXT: - .offset: 64
224 ; CHECK-NEXT: .size: 8
225 ; CHECK-NEXT: .value_kind: hidden_none
226 ; CHECK-NEXT: - .offset: 72
227 ; CHECK-NEXT: .size: 8
228 ; CHECK-NEXT: .value_kind: hidden_none
229 ; CHECK-NEXT: - .offset: 80
230 ; CHECK-NEXT: .size: 8
231 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
232 ; CHECK: .language: OpenCL C
233 ; CHECK-NEXT: .language_version:
234 ; CHECK-NEXT: - 2
235 ; CHECK-NEXT: - 0
236 ; CHECK: .name: test_ulong4
237 ; CHECK: .symbol: test_ulong4.kd
238 define amdgpu_kernel void @test_ulong4(<4 x i64> %a) #0
239 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !12
240 !kernel_arg_base_type !12 !kernel_arg_type_qual !4 {
241 ret void
242 }
244 ; CHECK: - .args:
245 ; CHECK-NEXT: - .name: a
246 ; CHECK-NEXT: .offset: 0
247 ; CHECK-NEXT: .size: 16
248 ; CHECK-NEXT: .type_name: half8
249 ; CHECK-NEXT: .value_kind: by_value
250 ; CHECK-NEXT: - .offset: 16
251 ; CHECK-NEXT: .size: 8
252 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
253 ; CHECK-NEXT: - .offset: 24
254 ; CHECK-NEXT: .size: 8
255 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
256 ; CHECK-NEXT: - .offset: 32
257 ; CHECK-NEXT: .size: 8
258 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
259 ; CHECK-NEXT: - .offset: 40
260 ; CHECK-NEXT: .size: 8
261 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
262 ; CHECK-NEXT: - .offset: 48
263 ; CHECK-NEXT: .size: 8
264 ; CHECK-NEXT: .value_kind: hidden_none
265 ; CHECK-NEXT: - .offset: 56
266 ; CHECK-NEXT: .size: 8
267 ; CHECK-NEXT: .value_kind: hidden_none
268 ; CHECK-NEXT: - .offset: 64
269 ; CHECK-NEXT: .size: 8
270 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
271 ; CHECK: .language: OpenCL C
272 ; CHECK-NEXT: .language_version:
273 ; CHECK-NEXT: - 2
274 ; CHECK-NEXT: - 0
275 ; CHECK: .name: test_half8
276 ; CHECK: .symbol: test_half8.kd
277 define amdgpu_kernel void @test_half8(<8 x half> %a) #0
278 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !13
279 !kernel_arg_base_type !13 !kernel_arg_type_qual !4 {
280 ret void
281 }
283 ; CHECK: - .args:
284 ; CHECK-NEXT: - .name: a
285 ; CHECK-NEXT: .offset: 0
286 ; CHECK-NEXT: .size: 64
287 ; CHECK-NEXT: .type_name: float16
288 ; CHECK-NEXT: .value_kind: by_value
289 ; CHECK-NEXT: - .offset: 64
290 ; CHECK-NEXT: .size: 8
291 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
292 ; CHECK-NEXT: - .offset: 72
293 ; CHECK-NEXT: .size: 8
294 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
295 ; CHECK-NEXT: - .offset: 80
296 ; CHECK-NEXT: .size: 8
297 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
298 ; CHECK-NEXT: - .offset: 88
299 ; CHECK-NEXT: .size: 8
300 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
301 ; CHECK-NEXT: - .offset: 96
302 ; CHECK-NEXT: .size: 8
303 ; CHECK-NEXT: .value_kind: hidden_none
304 ; CHECK-NEXT: - .offset: 104
305 ; CHECK-NEXT: .size: 8
306 ; CHECK-NEXT: .value_kind: hidden_none
307 ; CHECK-NEXT: - .offset: 112
308 ; CHECK-NEXT: .size: 8
309 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
310 ; CHECK: .language: OpenCL C
311 ; CHECK-NEXT: .language_version:
312 ; CHECK-NEXT: - 2
313 ; CHECK-NEXT: - 0
314 ; CHECK: .name: test_float16
315 ; CHECK: .symbol: test_float16.kd
316 define amdgpu_kernel void @test_float16(<16 x float> %a) #0
317 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !14
318 !kernel_arg_base_type !14 !kernel_arg_type_qual !4 {
319 ret void
320 }
322 ; CHECK: - .args:
323 ; CHECK-NEXT: - .name: a
324 ; CHECK-NEXT: .offset: 0
325 ; CHECK-NEXT: .size: 128
326 ; CHECK-NEXT: .type_name: double16
327 ; CHECK-NEXT: .value_kind: by_value
328 ; CHECK-NEXT: - .offset: 128
329 ; CHECK-NEXT: .size: 8
330 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
331 ; CHECK-NEXT: - .offset: 136
332 ; CHECK-NEXT: .size: 8
333 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
334 ; CHECK-NEXT: - .offset: 144
335 ; CHECK-NEXT: .size: 8
336 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
337 ; CHECK-NEXT: - .offset: 152
338 ; CHECK-NEXT: .size: 8
339 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
340 ; CHECK-NEXT: - .offset: 160
341 ; CHECK-NEXT: .size: 8
342 ; CHECK-NEXT: .value_kind: hidden_none
343 ; CHECK-NEXT: - .offset: 168
344 ; CHECK-NEXT: .size: 8
345 ; CHECK-NEXT: .value_kind: hidden_none
346 ; CHECK-NEXT: - .offset: 176
347 ; CHECK-NEXT: .size: 8
348 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
349 ; CHECK: .language: OpenCL C
350 ; CHECK-NEXT: .language_version:
351 ; CHECK-NEXT: - 2
352 ; CHECK-NEXT: - 0
353 ; CHECK: .name: test_double16
354 ; CHECK: .symbol: test_double16.kd
355 define amdgpu_kernel void @test_double16(<16 x double> %a) #0
356 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !15
357 !kernel_arg_base_type !15 !kernel_arg_type_qual !4 {
358 ret void
359 }
361 ; CHECK: - .args:
362 ; CHECK-NEXT: - .address_space: global
363 ; CHECK-NEXT: .name: a
364 ; CHECK-NEXT: .offset: 0
365 ; CHECK-NEXT: .size: 8
366 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
367 ; CHECK-NEXT: .value_kind: global_buffer
368 ; CHECK-NEXT: - .offset: 8
369 ; CHECK-NEXT: .size: 8
370 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
371 ; CHECK-NEXT: - .offset: 16
372 ; CHECK-NEXT: .size: 8
373 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
374 ; CHECK-NEXT: - .offset: 24
375 ; CHECK-NEXT: .size: 8
376 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
377 ; CHECK-NEXT: - .offset: 32
378 ; CHECK-NEXT: .size: 8
379 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
380 ; CHECK-NEXT: - .offset: 40
381 ; CHECK-NEXT: .size: 8
382 ; CHECK-NEXT: .value_kind: hidden_none
383 ; CHECK-NEXT: - .offset: 48
384 ; CHECK-NEXT: .size: 8
385 ; CHECK-NEXT: .value_kind: hidden_none
386 ; CHECK-NEXT: - .offset: 56
387 ; CHECK-NEXT: .size: 8
388 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
389 ; CHECK: .language: OpenCL C
390 ; CHECK-NEXT: .language_version:
391 ; CHECK-NEXT: - 2
392 ; CHECK-NEXT: - 0
393 ; CHECK: .name: test_pointer
394 ; CHECK: .symbol: test_pointer.kd
395 define amdgpu_kernel void @test_pointer(ptr addrspace(1) %a) #0
396 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !16
397 !kernel_arg_base_type !16 !kernel_arg_type_qual !4 {
398 ret void
399 }
401 ; CHECK: - .args:
402 ; CHECK-NEXT: - .name: a
403 ; CHECK-NEXT: .offset: 0
404 ; CHECK-NEXT: .size: 8
405 ; CHECK-NEXT: .type_name: image2d_t
406 ; CHECK-NEXT: .value_kind: image
407 ; CHECK-NEXT: - .offset: 8
408 ; CHECK-NEXT: .size: 8
409 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
410 ; CHECK-NEXT: - .offset: 16
411 ; CHECK-NEXT: .size: 8
412 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
413 ; CHECK-NEXT: - .offset: 24
414 ; CHECK-NEXT: .size: 8
415 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
416 ; CHECK-NEXT: - .offset: 32
417 ; CHECK-NEXT: .size: 8
418 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
419 ; CHECK-NEXT: - .offset: 40
420 ; CHECK-NEXT: .size: 8
421 ; CHECK-NEXT: .value_kind: hidden_none
422 ; CHECK-NEXT: - .offset: 48
423 ; CHECK-NEXT: .size: 8
424 ; CHECK-NEXT: .value_kind: hidden_none
425 ; CHECK-NEXT: - .offset: 56
426 ; CHECK-NEXT: .size: 8
427 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
428 ; CHECK: .language: OpenCL C
429 ; CHECK-NEXT: .language_version:
430 ; CHECK-NEXT: - 2
431 ; CHECK-NEXT: - 0
432 ; CHECK: .name: test_image
433 ; CHECK: .symbol: test_image.kd
434 define amdgpu_kernel void @test_image(ptr addrspace(1) %a) #0
435 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !17
436 !kernel_arg_base_type !17 !kernel_arg_type_qual !4 {
437 ret void
438 }
440 ; CHECK: - .args:
441 ; CHECK-NEXT: - .name: a
442 ; CHECK-NEXT: .offset: 0
443 ; CHECK-NEXT: .size: 4
444 ; CHECK-NEXT: .type_name: sampler_t
445 ; CHECK-NEXT: .value_kind: sampler
446 ; CHECK-NEXT: - .offset: 8
447 ; CHECK-NEXT: .size: 8
448 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
449 ; CHECK-NEXT: - .offset: 16
450 ; CHECK-NEXT: .size: 8
451 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
452 ; CHECK-NEXT: - .offset: 24
453 ; CHECK-NEXT: .size: 8
454 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
455 ; CHECK-NEXT: - .offset: 32
456 ; CHECK-NEXT: .size: 8
457 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
458 ; CHECK-NEXT: - .offset: 40
459 ; CHECK-NEXT: .size: 8
460 ; CHECK-NEXT: .value_kind: hidden_none
461 ; CHECK-NEXT: - .offset: 48
462 ; CHECK-NEXT: .size: 8
463 ; CHECK-NEXT: .value_kind: hidden_none
464 ; CHECK-NEXT: - .offset: 56
465 ; CHECK-NEXT: .size: 8
466 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
467 ; CHECK: .language: OpenCL C
468 ; CHECK-NEXT: .language_version:
469 ; CHECK-NEXT: - 2
470 ; CHECK-NEXT: - 0
471 ; CHECK: .name: test_sampler
472 ; CHECK: .symbol: test_sampler.kd
473 define amdgpu_kernel void @test_sampler(i32 %a) #0
474 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !18
475 !kernel_arg_base_type !18 !kernel_arg_type_qual !4 {
476 ret void
477 }
479 ; CHECK: - .args:
480 ; CHECK-NEXT: - .name: a
481 ; CHECK-NEXT: .offset: 0
482 ; CHECK-NEXT: .size: 8
483 ; CHECK-NEXT: .type_name: queue_t
484 ; CHECK-NEXT: .value_kind: queue
485 ; CHECK-NEXT: - .offset: 8
486 ; CHECK-NEXT: .size: 8
487 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
488 ; CHECK-NEXT: - .offset: 16
489 ; CHECK-NEXT: .size: 8
490 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
491 ; CHECK-NEXT: - .offset: 24
492 ; CHECK-NEXT: .size: 8
493 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
494 ; CHECK-NEXT: - .offset: 32
495 ; CHECK-NEXT: .size: 8
496 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
497 ; CHECK-NEXT: - .offset: 40
498 ; CHECK-NEXT: .size: 8
499 ; CHECK-NEXT: .value_kind: hidden_none
500 ; CHECK-NEXT: - .offset: 48
501 ; CHECK-NEXT: .size: 8
502 ; CHECK-NEXT: .value_kind: hidden_none
503 ; CHECK-NEXT: - .offset: 56
504 ; CHECK-NEXT: .size: 8
505 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
506 ; CHECK: .language: OpenCL C
507 ; CHECK-NEXT: .language_version:
508 ; CHECK-NEXT: - 2
509 ; CHECK-NEXT: - 0
510 ; CHECK: .name: test_queue
511 ; CHECK: .symbol: test_queue.kd
512 define amdgpu_kernel void @test_queue(ptr addrspace(1) %a) #0
513 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !19
514 !kernel_arg_base_type !19 !kernel_arg_type_qual !4 {
515 ret void
516 }
518 ; CHECK: - .args:
519 ; CHECK-NEXT: .name: a
520 ; CHECK-NEXT: .offset: 0
521 ; CHECK-NEXT: .size: 8
522 ; CHECK-NEXT: .type_name: struct A
523 ; CHECK-NEXT: .value_kind: by_value
524 ; CHECK-NEXT: - .offset: 8
525 ; CHECK-NEXT: .size: 8
526 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
527 ; CHECK-NEXT: - .offset: 16
528 ; CHECK-NEXT: .size: 8
529 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
530 ; CHECK-NEXT: - .offset: 24
531 ; CHECK-NEXT: .size: 8
532 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
533 ; CHECK-NEXT: - .offset: 32
534 ; CHECK-NEXT: .size: 8
535 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
536 ; CHECK-NEXT: - .offset: 40
537 ; CHECK-NEXT: .size: 8
538 ; CHECK-NEXT: .value_kind: hidden_none
539 ; CHECK-NEXT: - .offset: 48
540 ; CHECK-NEXT: .size: 8
541 ; CHECK-NEXT: .value_kind: hidden_none
542 ; CHECK-NEXT: - .offset: 56
543 ; CHECK-NEXT: .size: 8
544 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
545 ; CHECK: .language: OpenCL C
546 ; CHECK-NEXT: .language_version:
547 ; CHECK-NEXT: - 2
548 ; CHECK-NEXT: - 0
549 ; CHECK: .name: test_struct
550 ; CHECK: .symbol: test_struct.kd
551 define amdgpu_kernel void @test_struct(%struct.A %a) #0
552 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20
553 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 {
554 ret void
555 }
557 ; CHECK: - .args:
558 ; CHECK-NEXT: .name: a
559 ; CHECK-NEXT: .offset: 0
560 ; CHECK-NEXT: .size: 8
561 ; CHECK-NEXT: .type_name: struct A
562 ; CHECK-NEXT: .value_kind: by_value
563 ; CHECK-NEXT: - .offset: 8
564 ; CHECK-NEXT: .size: 8
565 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
566 ; CHECK-NEXT: - .offset: 16
567 ; CHECK-NEXT: .size: 8
568 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
569 ; CHECK-NEXT: - .offset: 24
570 ; CHECK-NEXT: .size: 8
571 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
572 ; CHECK-NEXT: - .offset: 32
573 ; CHECK-NEXT: .size: 8
574 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
575 ; CHECK-NEXT: - .offset: 40
576 ; CHECK-NEXT: .size: 8
577 ; CHECK-NEXT: .value_kind: hidden_none
578 ; CHECK-NEXT: - .offset: 48
579 ; CHECK-NEXT: .size: 8
580 ; CHECK-NEXT: .value_kind: hidden_none
581 ; CHECK-NEXT: - .offset: 56
582 ; CHECK-NEXT: .size: 8
583 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
584 ; CHECK: .language: OpenCL C
585 ; CHECK-NEXT: .language_version:
586 ; CHECK-NEXT: - 2
587 ; CHECK-NEXT: - 0
588 ; CHECK: .name: test_struct_byref_constant
589 ; CHECK: .symbol: test_struct_byref_constant.kd
590 define amdgpu_kernel void @test_struct_byref_constant(ptr addrspace(4) byref(%struct.A) %a) #0
591 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20
592 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 {
593 ret void
594 }
596 ; CHECK: - .args:
597 ; CHECK-NEXT: .name: a
598 ; CHECK-NEXT: .offset: 0
599 ; CHECK-NEXT: .size: 32
600 ; CHECK-NEXT: .type_name: struct A
601 ; CHECK-NEXT: .value_kind: by_value
602 ; CHECK-NEXT: - .offset: 32
603 ; CHECK-NEXT: .size: 8
604 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
605 ; CHECK-NEXT: - .offset: 40
606 ; CHECK-NEXT: .size: 8
607 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
608 ; CHECK-NEXT: - .offset: 48
609 ; CHECK-NEXT: .size: 8
610 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
611 ; CHECK-NEXT: - .offset: 56
612 ; CHECK-NEXT: .size: 8
613 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
614 ; CHECK-NEXT: - .offset: 64
615 ; CHECK-NEXT: .size: 8
616 ; CHECK-NEXT: .value_kind: hidden_none
617 ; CHECK-NEXT: - .offset: 72
618 ; CHECK-NEXT: .size: 8
619 ; CHECK-NEXT: .value_kind: hidden_none
620 ; CHECK-NEXT: - .offset: 80
621 ; CHECK-NEXT: .size: 8
622 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
623 ; CHECK: .language: OpenCL C
624 ; CHECK-NEXT: .language_version:
625 ; CHECK-NEXT: - 2
626 ; CHECK-NEXT: - 0
627 ; CHECK: .name: test_array
628 ; CHECK: .symbol: test_array.kd
629 define amdgpu_kernel void @test_array([32 x i8] %a) #0
630 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20
631 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 {
632 ret void
633 }
635 ; CHECK: - .args:
636 ; CHECK-NEXT: .name: a
637 ; CHECK-NEXT: .offset: 0
638 ; CHECK-NEXT: .size: 32
639 ; CHECK-NEXT: .type_name: struct A
640 ; CHECK-NEXT: .value_kind: by_value
641 ; CHECK-NEXT: - .offset: 32
642 ; CHECK-NEXT: .size: 8
643 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
644 ; CHECK-NEXT: - .offset: 40
645 ; CHECK-NEXT: .size: 8
646 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
647 ; CHECK-NEXT: - .offset: 48
648 ; CHECK-NEXT: .size: 8
649 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
650 ; CHECK-NEXT: - .offset: 56
651 ; CHECK-NEXT: .size: 8
652 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
653 ; CHECK-NEXT: - .offset: 64
654 ; CHECK-NEXT: .size: 8
655 ; CHECK-NEXT: .value_kind: hidden_none
656 ; CHECK-NEXT: - .offset: 72
657 ; CHECK-NEXT: .size: 8
658 ; CHECK-NEXT: .value_kind: hidden_none
659 ; CHECK-NEXT: - .offset: 80
660 ; CHECK-NEXT: .size: 8
661 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
662 ; CHECK: .language: OpenCL C
663 ; CHECK-NEXT: .language_version:
664 ; CHECK-NEXT: - 2
665 ; CHECK-NEXT: - 0
666 ; CHECK: .name: test_array_byref_constant
667 ; CHECK: .symbol: test_array_byref_constant.kd
668 define amdgpu_kernel void @test_array_byref_constant(ptr addrspace(4) byref([32 x i8]) %a) #0
669 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20
670 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 {
671 ret void
672 }
674 ; CHECK: - .args:
675 ; CHECK-NEXT: - .name: a
676 ; CHECK-NEXT: .offset: 0
677 ; CHECK-NEXT: .size: 16
678 ; CHECK-NEXT: .type_name: i128
679 ; CHECK-NEXT: .value_kind: by_value
680 ; CHECK-NEXT: - .offset: 16
681 ; CHECK-NEXT: .size: 8
682 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
683 ; CHECK-NEXT: - .offset: 24
684 ; CHECK-NEXT: .size: 8
685 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
686 ; CHECK-NEXT: - .offset: 32
687 ; CHECK-NEXT: .size: 8
688 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
689 ; CHECK-NEXT: - .offset: 40
690 ; CHECK-NEXT: .size: 8
691 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
692 ; CHECK-NEXT: - .offset: 48
693 ; CHECK-NEXT: .size: 8
694 ; CHECK-NEXT: .value_kind: hidden_none
695 ; CHECK-NEXT: - .offset: 56
696 ; CHECK-NEXT: .size: 8
697 ; CHECK-NEXT: .value_kind: hidden_none
698 ; CHECK-NEXT: - .offset: 64
699 ; CHECK-NEXT: .size: 8
700 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
701 ; CHECK: .language: OpenCL C
702 ; CHECK-NEXT: .language_version:
703 ; CHECK-NEXT: - 2
704 ; CHECK-NEXT: - 0
705 ; CHECK: .name: test_i128
706 ; CHECK: .symbol: test_i128.kd
707 define amdgpu_kernel void @test_i128(i128 %a) #0
708 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !21
709 !kernel_arg_base_type !21 !kernel_arg_type_qual !4 {
710 ret void
711 }
713 ; CHECK: - .args:
714 ; CHECK-NEXT: - .name: a
715 ; CHECK-NEXT: .offset: 0
716 ; CHECK-NEXT: .size: 4
717 ; CHECK-NEXT: .type_name: int
718 ; CHECK-NEXT: .value_kind: by_value
719 ; CHECK-NEXT: - .name: b
720 ; CHECK-NEXT: .offset: 4
721 ; CHECK-NEXT: .size: 4
722 ; CHECK-NEXT: .type_name: short2
723 ; CHECK-NEXT: .value_kind: by_value
724 ; CHECK-NEXT: - .name: c
725 ; CHECK-NEXT: .offset: 8
726 ; CHECK-NEXT: .size: 4
727 ; CHECK-NEXT: .type_name: char3
728 ; CHECK-NEXT: .value_kind: by_value
729 ; CHECK-NEXT: - .offset: 16
730 ; CHECK-NEXT: .size: 8
731 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
732 ; CHECK-NEXT: - .offset: 24
733 ; CHECK-NEXT: .size: 8
734 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
735 ; CHECK-NEXT: - .offset: 32
736 ; CHECK-NEXT: .size: 8
737 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
738 ; CHECK-NEXT: - .offset: 40
739 ; CHECK-NEXT: .size: 8
740 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
741 ; CHECK-NEXT: - .offset: 48
742 ; CHECK-NEXT: .size: 8
743 ; CHECK-NEXT: .value_kind: hidden_none
744 ; CHECK-NEXT: - .offset: 56
745 ; CHECK-NEXT: .size: 8
746 ; CHECK-NEXT: .value_kind: hidden_none
747 ; CHECK-NEXT: - .offset: 64
748 ; CHECK-NEXT: .size: 8
749 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
750 ; CHECK: .language: OpenCL C
751 ; CHECK-NEXT: .language_version:
752 ; CHECK-NEXT: - 2
753 ; CHECK-NEXT: - 0
754 ; CHECK: .name: test_multi_arg
755 ; CHECK: .symbol: test_multi_arg.kd
756 define amdgpu_kernel void @test_multi_arg(i32 %a, <2 x i16> %b, <3 x i8> %c) #0
757 !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !24
758 !kernel_arg_base_type !24 !kernel_arg_type_qual !25 {
759 ret void
760 }
762 ; CHECK: - .args:
763 ; CHECK-NEXT: - .address_space: global
764 ; CHECK-NEXT: .name: g
765 ; CHECK-NEXT: .offset: 0
766 ; CHECK-NEXT: .size: 8
767 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
768 ; CHECK-NEXT: .value_kind: global_buffer
769 ; CHECK-NEXT: - .address_space: constant
770 ; CHECK-NEXT: .name: c
771 ; CHECK-NEXT: .offset: 8
772 ; CHECK-NEXT: .size: 8
773 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
774 ; CHECK-NEXT: .value_kind: global_buffer
775 ; CHECK-NEXT: - .address_space: local
776 ; CHECK-NEXT: .name: l
777 ; CHECK-NEXT: .offset: 16
778 ; CHECK-NEXT: .pointee_align: 4
779 ; CHECK-NEXT: .size: 4
780 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
781 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
782 ; CHECK-NEXT: - .offset: 24
783 ; CHECK-NEXT: .size: 8
784 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
785 ; CHECK-NEXT: - .offset: 32
786 ; CHECK-NEXT: .size: 8
787 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
788 ; CHECK-NEXT: - .offset: 40
789 ; CHECK-NEXT: .size: 8
790 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
791 ; CHECK-NEXT: - .offset: 48
792 ; CHECK-NEXT: .size: 8
793 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
794 ; CHECK-NEXT: - .offset: 56
795 ; CHECK-NEXT: .size: 8
796 ; CHECK-NEXT: .value_kind: hidden_none
797 ; CHECK-NEXT: - .offset: 64
798 ; CHECK-NEXT: .size: 8
799 ; CHECK-NEXT: .value_kind: hidden_none
800 ; CHECK-NEXT: - .offset: 72
801 ; CHECK-NEXT: .size: 8
802 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
803 ; CHECK: .language: OpenCL C
804 ; CHECK-NEXT: .language_version:
805 ; CHECK-NEXT: - 2
806 ; CHECK-NEXT: - 0
807 ; CHECK: .name: test_addr_space
808 ; CHECK: .symbol: test_addr_space.kd
809 define amdgpu_kernel void @test_addr_space(ptr addrspace(1) %g,
810 ptr addrspace(4) %c,
811 ptr addrspace(3) align 4 %l) #0
812 !kernel_arg_addr_space !50 !kernel_arg_access_qual !23 !kernel_arg_type !51
813 !kernel_arg_base_type !51 !kernel_arg_type_qual !25 {
814 ret void
815 }
817 ; CHECK: - .args:
818 ; CHECK-NEXT: - .address_space: global
819 ; CHECK-NEXT: .is_volatile: true
820 ; CHECK-NEXT: .name: a
821 ; CHECK-NEXT: .offset: 0
822 ; CHECK-NEXT: .size: 8
823 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
824 ; CHECK-NEXT: .value_kind: global_buffer
825 ; CHECK-NEXT: - .address_space: global
826 ; CHECK-NEXT: .is_const: true
827 ; CHECK-NEXT: .is_restrict: true
828 ; CHECK-NEXT: .name: b
829 ; CHECK-NEXT: .offset: 8
830 ; CHECK-NEXT: .size: 8
831 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
832 ; CHECK-NEXT: .value_kind: global_buffer
833 ; CHECK-NEXT: - .is_pipe: true
834 ; CHECK-NEXT: .name: c
835 ; CHECK-NEXT: .offset: 16
836 ; CHECK-NEXT: .size: 8
837 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
838 ; CHECK-NEXT: .value_kind: pipe
839 ; CHECK-NEXT: - .offset: 24
840 ; CHECK-NEXT: .size: 8
841 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
842 ; CHECK-NEXT: - .offset: 32
843 ; CHECK-NEXT: .size: 8
844 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
845 ; CHECK-NEXT: - .offset: 40
846 ; CHECK-NEXT: .size: 8
847 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
848 ; CHECK-NEXT: - .offset: 48
849 ; CHECK-NEXT: .size: 8
850 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
851 ; CHECK-NEXT: - .offset: 56
852 ; CHECK-NEXT: .size: 8
853 ; CHECK-NEXT: .value_kind: hidden_none
854 ; CHECK-NEXT: - .offset: 64
855 ; CHECK-NEXT: .size: 8
856 ; CHECK-NEXT: .value_kind: hidden_none
857 ; CHECK-NEXT: - .offset: 72
858 ; CHECK-NEXT: .size: 8
859 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
860 ; CHECK: .language: OpenCL C
861 ; CHECK-NEXT: .language_version:
862 ; CHECK-NEXT: - 2
863 ; CHECK-NEXT: - 0
864 ; CHECK: .name: test_type_qual
865 ; CHECK: .symbol: test_type_qual.kd
866 define amdgpu_kernel void @test_type_qual(ptr addrspace(1) %a,
867 ptr addrspace(1) %b,
868 ptr addrspace(1) %c) #0
869 !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !51
870 !kernel_arg_base_type !51 !kernel_arg_type_qual !70 {
871 ret void
872 }
874 ; CHECK: - .args:
875 ; CHECK-NEXT: - .access: read_only
876 ; CHECK-NEXT: .name: ro
877 ; CHECK-NEXT: .offset: 0
878 ; CHECK-NEXT: .size: 8
879 ; CHECK-NEXT: .type_name: image1d_t
880 ; CHECK-NEXT: .value_kind: image
881 ; CHECK-NEXT: - .access: write_only
882 ; CHECK-NEXT: .name: wo
883 ; CHECK-NEXT: .offset: 8
884 ; CHECK-NEXT: .size: 8
885 ; CHECK-NEXT: .type_name: image2d_t
886 ; CHECK-NEXT: .value_kind: image
887 ; CHECK-NEXT: - .access: read_write
888 ; CHECK-NEXT: .name: rw
889 ; CHECK-NEXT: .offset: 16
890 ; CHECK-NEXT: .size: 8
891 ; CHECK-NEXT: .type_name: image3d_t
892 ; CHECK-NEXT: .value_kind: image
893 ; CHECK-NEXT: - .offset: 24
894 ; CHECK-NEXT: .size: 8
895 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
896 ; CHECK-NEXT: - .offset: 32
897 ; CHECK-NEXT: .size: 8
898 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
899 ; CHECK-NEXT: - .offset: 40
900 ; CHECK-NEXT: .size: 8
901 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
902 ; CHECK-NEXT: - .offset: 48
903 ; CHECK-NEXT: .size: 8
904 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
905 ; CHECK-NEXT: - .offset: 56
906 ; CHECK-NEXT: .size: 8
907 ; CHECK-NEXT: .value_kind: hidden_none
908 ; CHECK-NEXT: - .offset: 64
909 ; CHECK-NEXT: .size: 8
910 ; CHECK-NEXT: .value_kind: hidden_none
911 ; CHECK-NEXT: - .offset: 72
912 ; CHECK-NEXT: .size: 8
913 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
914 ; CHECK: .language: OpenCL C
915 ; CHECK-NEXT: .language_version:
916 ; CHECK-NEXT: - 2
917 ; CHECK-NEXT: - 0
918 ; CHECK: .name: test_access_qual
919 ; CHECK: .symbol: test_access_qual.kd
920 define amdgpu_kernel void @test_access_qual(ptr addrspace(1) %ro,
921 ptr addrspace(1) %wo,
922 ptr addrspace(1) %rw) #0
923 !kernel_arg_addr_space !60 !kernel_arg_access_qual !61 !kernel_arg_type !62
924 !kernel_arg_base_type !62 !kernel_arg_type_qual !25 {
925 ret void
926 }
928 ; CHECK: - .args:
929 ; CHECK-NEXT: - .name: a
930 ; CHECK-NEXT: .offset: 0
931 ; CHECK-NEXT: .size: 4
932 ; CHECK-NEXT: .type_name: int
933 ; CHECK-NEXT: .value_kind: by_value
934 ; CHECK-NEXT: - .offset: 8
935 ; CHECK-NEXT: .size: 8
936 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
937 ; CHECK-NEXT: - .offset: 16
938 ; CHECK-NEXT: .size: 8
939 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
940 ; CHECK-NEXT: - .offset: 24
941 ; CHECK-NEXT: .size: 8
942 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
943 ; CHECK-NEXT: - .offset: 32
944 ; CHECK-NEXT: .size: 8
945 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
946 ; CHECK-NEXT: - .offset: 40
947 ; CHECK-NEXT: .size: 8
948 ; CHECK-NEXT: .value_kind: hidden_none
949 ; CHECK-NEXT: - .offset: 48
950 ; CHECK-NEXT: .size: 8
951 ; CHECK-NEXT: .value_kind: hidden_none
952 ; CHECK-NEXT: - .offset: 56
953 ; CHECK-NEXT: .size: 8
954 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
955 ; CHECK: .language: OpenCL C
956 ; CHECK-NEXT: .language_version:
957 ; CHECK-NEXT: - 2
958 ; CHECK-NEXT: - 0
959 ; CHECK: .name: test_vec_type_hint_half
960 ; CHECK: .symbol: test_vec_type_hint_half.kd
961 ; CHECK: .vec_type_hint: half
962 define amdgpu_kernel void @test_vec_type_hint_half(i32 %a) #0
963 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
964 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !26 {
965 ret void
966 }
968 ; CHECK: - .args:
969 ; CHECK-NEXT: - .name: a
970 ; CHECK-NEXT: .offset: 0
971 ; CHECK-NEXT: .size: 4
972 ; CHECK-NEXT: .type_name: int
973 ; CHECK-NEXT: .value_kind: by_value
974 ; CHECK-NEXT: - .offset: 8
975 ; CHECK-NEXT: .size: 8
976 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
977 ; CHECK-NEXT: - .offset: 16
978 ; CHECK-NEXT: .size: 8
979 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
980 ; CHECK-NEXT: - .offset: 24
981 ; CHECK-NEXT: .size: 8
982 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
983 ; CHECK-NEXT: - .offset: 32
984 ; CHECK-NEXT: .size: 8
985 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
986 ; CHECK-NEXT: - .offset: 40
987 ; CHECK-NEXT: .size: 8
988 ; CHECK-NEXT: .value_kind: hidden_none
989 ; CHECK-NEXT: - .offset: 48
990 ; CHECK-NEXT: .size: 8
991 ; CHECK-NEXT: .value_kind: hidden_none
992 ; CHECK-NEXT: - .offset: 56
993 ; CHECK-NEXT: .size: 8
994 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
995 ; CHECK: .language: OpenCL C
996 ; CHECK-NEXT: .language_version:
997 ; CHECK-NEXT: - 2
998 ; CHECK-NEXT: - 0
999 ; CHECK: .name: test_vec_type_hint_float
1000 ; CHECK: .symbol: test_vec_type_hint_float.kd
1001 ; CHECK: .vec_type_hint: float
1002 define amdgpu_kernel void @test_vec_type_hint_float(i32 %a) #0
1003 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
1004 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !27 {
1005 ret void
1006 }
1008 ; CHECK: - .args:
1009 ; CHECK-NEXT: - .name: a
1010 ; CHECK-NEXT: .offset: 0
1011 ; CHECK-NEXT: .size: 4
1012 ; CHECK-NEXT: .type_name: int
1013 ; CHECK-NEXT: .value_kind: by_value
1014 ; CHECK-NEXT: - .offset: 8
1015 ; CHECK-NEXT: .size: 8
1016 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1017 ; CHECK-NEXT: - .offset: 16
1018 ; CHECK-NEXT: .size: 8
1019 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1020 ; CHECK-NEXT: - .offset: 24
1021 ; CHECK-NEXT: .size: 8
1022 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1023 ; CHECK-NEXT: - .offset: 32
1024 ; CHECK-NEXT: .size: 8
1025 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1026 ; CHECK-NEXT: - .offset: 40
1027 ; CHECK-NEXT: .size: 8
1028 ; CHECK-NEXT: .value_kind: hidden_none
1029 ; CHECK-NEXT: - .offset: 48
1030 ; CHECK-NEXT: .size: 8
1031 ; CHECK-NEXT: .value_kind: hidden_none
1032 ; CHECK-NEXT: - .offset: 56
1033 ; CHECK-NEXT: .size: 8
1034 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1035 ; CHECK: .language: OpenCL C
1036 ; CHECK-NEXT: .language_version:
1037 ; CHECK-NEXT: - 2
1038 ; CHECK-NEXT: - 0
1039 ; CHECK: .name: test_vec_type_hint_double
1040 ; CHECK: .symbol: test_vec_type_hint_double.kd
1041 ; CHECK: .vec_type_hint: double
1042 define amdgpu_kernel void @test_vec_type_hint_double(i32 %a) #0
1043 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
1044 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !28 {
1045 ret void
1046 }
1048 ; CHECK: - .args:
1049 ; CHECK-NEXT: - .name: a
1050 ; CHECK-NEXT: .offset: 0
1051 ; CHECK-NEXT: .size: 4
1052 ; CHECK-NEXT: .type_name: int
1053 ; CHECK-NEXT: .value_kind: by_value
1054 ; CHECK-NEXT: - .offset: 8
1055 ; CHECK-NEXT: .size: 8
1056 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1057 ; CHECK-NEXT: - .offset: 16
1058 ; CHECK-NEXT: .size: 8
1059 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1060 ; CHECK-NEXT: - .offset: 24
1061 ; CHECK-NEXT: .size: 8
1062 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1063 ; CHECK-NEXT: - .offset: 32
1064 ; CHECK-NEXT: .size: 8
1065 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1066 ; CHECK-NEXT: - .offset: 40
1067 ; CHECK-NEXT: .size: 8
1068 ; CHECK-NEXT: .value_kind: hidden_none
1069 ; CHECK-NEXT: - .offset: 48
1070 ; CHECK-NEXT: .size: 8
1071 ; CHECK-NEXT: .value_kind: hidden_none
1072 ; CHECK-NEXT: - .offset: 56
1073 ; CHECK-NEXT: .size: 8
1074 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1075 ; CHECK: .language: OpenCL C
1076 ; CHECK-NEXT: .language_version:
1077 ; CHECK-NEXT: - 2
1078 ; CHECK-NEXT: - 0
1079 ; CHECK: .name: test_vec_type_hint_char
1080 ; CHECK: .symbol: test_vec_type_hint_char.kd
1081 ; CHECK: .vec_type_hint: char
1082 define amdgpu_kernel void @test_vec_type_hint_char(i32 %a) #0
1083 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
1084 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !29 {
1085 ret void
1086 }
1088 ; CHECK: - .args:
1089 ; CHECK-NEXT: - .name: a
1090 ; CHECK-NEXT: .offset: 0
1091 ; CHECK-NEXT: .size: 4
1092 ; CHECK-NEXT: .type_name: int
1093 ; CHECK-NEXT: .value_kind: by_value
1094 ; CHECK-NEXT: - .offset: 8
1095 ; CHECK-NEXT: .size: 8
1096 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1097 ; CHECK-NEXT: - .offset: 16
1098 ; CHECK-NEXT: .size: 8
1099 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1100 ; CHECK-NEXT: - .offset: 24
1101 ; CHECK-NEXT: .size: 8
1102 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1103 ; CHECK-NEXT: - .offset: 32
1104 ; CHECK-NEXT: .size: 8
1105 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1106 ; CHECK-NEXT: - .offset: 40
1107 ; CHECK-NEXT: .size: 8
1108 ; CHECK-NEXT: .value_kind: hidden_none
1109 ; CHECK-NEXT: - .offset: 48
1110 ; CHECK-NEXT: .size: 8
1111 ; CHECK-NEXT: .value_kind: hidden_none
1112 ; CHECK-NEXT: - .offset: 56
1113 ; CHECK-NEXT: .size: 8
1114 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1115 ; CHECK: .language: OpenCL C
1116 ; CHECK-NEXT: .language_version:
1117 ; CHECK-NEXT: - 2
1118 ; CHECK-NEXT: - 0
1119 ; CHECK: .name: test_vec_type_hint_short
1120 ; CHECK: .symbol: test_vec_type_hint_short.kd
1121 ; CHECK: .vec_type_hint: short
1122 define amdgpu_kernel void @test_vec_type_hint_short(i32 %a) #0
1123 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
1124 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !30 {
1125 ret void
1126 }
1128 ; CHECK: - .args:
1129 ; CHECK-NEXT: - .name: a
1130 ; CHECK-NEXT: .offset: 0
1131 ; CHECK-NEXT: .size: 4
1132 ; CHECK-NEXT: .type_name: int
1133 ; CHECK-NEXT: .value_kind: by_value
1134 ; CHECK-NEXT: - .offset: 8
1135 ; CHECK-NEXT: .size: 8
1136 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1137 ; CHECK-NEXT: - .offset: 16
1138 ; CHECK-NEXT: .size: 8
1139 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1140 ; CHECK-NEXT: - .offset: 24
1141 ; CHECK-NEXT: .size: 8
1142 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1143 ; CHECK-NEXT: - .offset: 32
1144 ; CHECK-NEXT: .size: 8
1145 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1146 ; CHECK-NEXT: - .offset: 40
1147 ; CHECK-NEXT: .size: 8
1148 ; CHECK-NEXT: .value_kind: hidden_none
1149 ; CHECK-NEXT: - .offset: 48
1150 ; CHECK-NEXT: .size: 8
1151 ; CHECK-NEXT: .value_kind: hidden_none
1152 ; CHECK-NEXT: - .offset: 56
1153 ; CHECK-NEXT: .size: 8
1154 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1155 ; CHECK: .language: OpenCL C
1156 ; CHECK-NEXT: .language_version:
1157 ; CHECK-NEXT: - 2
1158 ; CHECK-NEXT: - 0
1159 ; CHECK: .name: test_vec_type_hint_long
1160 ; CHECK: .symbol: test_vec_type_hint_long.kd
1161 ; CHECK: .vec_type_hint: long
1162 define amdgpu_kernel void @test_vec_type_hint_long(i32 %a) #0
1163 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
1164 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !31 {
1165 ret void
1166 }
1168 ; CHECK: - .args:
1169 ; CHECK-NEXT: - .name: a
1170 ; CHECK-NEXT: .offset: 0
1171 ; CHECK-NEXT: .size: 4
1172 ; CHECK-NEXT: .type_name: int
1173 ; CHECK-NEXT: .value_kind: by_value
1174 ; CHECK-NEXT: - .offset: 8
1175 ; CHECK-NEXT: .size: 8
1176 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1177 ; CHECK-NEXT: - .offset: 16
1178 ; CHECK-NEXT: .size: 8
1179 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1180 ; CHECK-NEXT: - .offset: 24
1181 ; CHECK-NEXT: .size: 8
1182 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1183 ; CHECK-NEXT: - .offset: 32
1184 ; CHECK-NEXT: .size: 8
1185 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1186 ; CHECK-NEXT: - .offset: 40
1187 ; CHECK-NEXT: .size: 8
1188 ; CHECK-NEXT: .value_kind: hidden_none
1189 ; CHECK-NEXT: - .offset: 48
1190 ; CHECK-NEXT: .size: 8
1191 ; CHECK-NEXT: .value_kind: hidden_none
1192 ; CHECK-NEXT: - .offset: 56
1193 ; CHECK-NEXT: .size: 8
1194 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1195 ; CHECK: .language: OpenCL C
1196 ; CHECK-NEXT: .language_version:
1197 ; CHECK-NEXT: - 2
1198 ; CHECK-NEXT: - 0
1199 ; CHECK: .name: test_vec_type_hint_unknown
1200 ; CHECK: .symbol: test_vec_type_hint_unknown.kd
1201 ; CHECK: .vec_type_hint: unknown
1202 define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %a) #0
1203 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
1204 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !32 {
1205 ret void
1206 }
1208 ; CHECK: - .args:
1209 ; CHECK-NEXT: - .name: a
1210 ; CHECK-NEXT: .offset: 0
1211 ; CHECK-NEXT: .size: 4
1212 ; CHECK-NEXT: .type_name: int
1213 ; CHECK-NEXT: .value_kind: by_value
1214 ; CHECK-NEXT: - .offset: 8
1215 ; CHECK-NEXT: .size: 8
1216 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1217 ; CHECK-NEXT: - .offset: 16
1218 ; CHECK-NEXT: .size: 8
1219 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1220 ; CHECK-NEXT: - .offset: 24
1221 ; CHECK-NEXT: .size: 8
1222 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1223 ; CHECK-NEXT: - .offset: 32
1224 ; CHECK-NEXT: .size: 8
1225 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1226 ; CHECK-NEXT: - .offset: 40
1227 ; CHECK-NEXT: .size: 8
1228 ; CHECK-NEXT: .value_kind: hidden_none
1229 ; CHECK-NEXT: - .offset: 48
1230 ; CHECK-NEXT: .size: 8
1231 ; CHECK-NEXT: .value_kind: hidden_none
1232 ; CHECK-NEXT: - .offset: 56
1233 ; CHECK-NEXT: .size: 8
1234 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1235 ; CHECK: .language: OpenCL C
1236 ; CHECK-NEXT: .language_version:
1237 ; CHECK-NEXT: - 2
1238 ; CHECK-NEXT: - 0
1239 ; CHECK: .name: test_reqd_wgs_vec_type_hint
1240 ; CHECK: .reqd_workgroup_size:
1241 ; CHECK-NEXT: - 1
1242 ; CHECK-NEXT: - 2
1243 ; CHECK-NEXT: - 4
1244 ; CHECK: .symbol: test_reqd_wgs_vec_type_hint.kd
1245 ; CHECK: .vec_type_hint: int
1246 define amdgpu_kernel void @test_reqd_wgs_vec_type_hint(i32 %a) #0
1247 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
1248 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !5
1249 !reqd_work_group_size !6 {
1250 ret void
1251 }
1253 ; CHECK: - .args:
1254 ; CHECK-NEXT: - .name: a
1255 ; CHECK-NEXT: .offset: 0
1256 ; CHECK-NEXT: .size: 4
1257 ; CHECK-NEXT: .type_name: int
1258 ; CHECK-NEXT: .value_kind: by_value
1259 ; CHECK-NEXT: - .offset: 8
1260 ; CHECK-NEXT: .size: 8
1261 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1262 ; CHECK-NEXT: - .offset: 16
1263 ; CHECK-NEXT: .size: 8
1264 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1265 ; CHECK-NEXT: - .offset: 24
1266 ; CHECK-NEXT: .size: 8
1267 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1268 ; CHECK-NEXT: - .offset: 32
1269 ; CHECK-NEXT: .size: 8
1270 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1271 ; CHECK-NEXT: - .offset: 40
1272 ; CHECK-NEXT: .size: 8
1273 ; CHECK-NEXT: .value_kind: hidden_none
1274 ; CHECK-NEXT: - .offset: 48
1275 ; CHECK-NEXT: .size: 8
1276 ; CHECK-NEXT: .value_kind: hidden_none
1277 ; CHECK-NEXT: - .offset: 56
1278 ; CHECK-NEXT: .size: 8
1279 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1280 ; CHECK: .language: OpenCL C
1281 ; CHECK-NEXT: .language_version:
1282 ; CHECK-NEXT: - 2
1283 ; CHECK-NEXT: - 0
1284 ; CHECK: .name: test_wgs_hint_vec_type_hint
1285 ; CHECK: .symbol: test_wgs_hint_vec_type_hint.kd
1286 ; CHECK: .vec_type_hint: uint4
1287 ; CHECK: .workgroup_size_hint:
1288 ; CHECK-NEXT: - 8
1289 ; CHECK-NEXT: - 16
1290 ; CHECK-NEXT: - 32
1291 define amdgpu_kernel void @test_wgs_hint_vec_type_hint(i32 %a) #0
1292 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
1293 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !7
1294 !work_group_size_hint !8 {
1295 ret void
1296 }
1298 ; CHECK: - .args:
1299 ; CHECK-NEXT: - .address_space: global
1300 ; CHECK-NEXT: .name: a
1301 ; CHECK-NEXT: .offset: 0
1302 ; CHECK-NEXT: .size: 8
1303 ; CHECK-NEXT: .type_name: 'int addrspace(5)* addrspace(5)*'
1304 ; CHECK-NEXT: .value_kind: global_buffer
1305 ; CHECK-NEXT: - .offset: 8
1306 ; CHECK-NEXT: .size: 8
1307 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1308 ; CHECK-NEXT: - .offset: 16
1309 ; CHECK-NEXT: .size: 8
1310 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1311 ; CHECK-NEXT: - .offset: 24
1312 ; CHECK-NEXT: .size: 8
1313 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1314 ; CHECK-NEXT: - .offset: 32
1315 ; CHECK-NEXT: .size: 8
1316 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1317 ; CHECK-NEXT: - .offset: 40
1318 ; CHECK-NEXT: .size: 8
1319 ; CHECK-NEXT: .value_kind: hidden_none
1320 ; CHECK-NEXT: - .offset: 48
1321 ; CHECK-NEXT: .size: 8
1322 ; CHECK-NEXT: .value_kind: hidden_none
1323 ; CHECK-NEXT: - .offset: 56
1324 ; CHECK-NEXT: .size: 8
1325 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1326 ; CHECK: .language: OpenCL C
1327 ; CHECK-NEXT: .language_version:
1328 ; CHECK-NEXT: - 2
1329 ; CHECK-NEXT: - 0
1330 ; CHECK: .name: test_arg_ptr_to_ptr
1331 ; CHECK: .symbol: test_arg_ptr_to_ptr.kd
1332 define amdgpu_kernel void @test_arg_ptr_to_ptr(ptr addrspace(1) %a) #0
1333 !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !80
1334 !kernel_arg_base_type !80 !kernel_arg_type_qual !4 {
1335 ret void
1336 }
1338 ; CHECK: - .args:
1339 ; CHECK-NEXT: .name: a
1340 ; CHECK-NEXT: .offset: 0
1341 ; CHECK-NEXT: .size: 8
1342 ; CHECK-NEXT: .type_name: struct B
1343 ; CHECK-NEXT: .value_kind: by_value
1344 ; CHECK-NEXT: - .offset: 8
1345 ; CHECK-NEXT: .size: 8
1346 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1347 ; CHECK-NEXT: - .offset: 16
1348 ; CHECK-NEXT: .size: 8
1349 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1350 ; CHECK-NEXT: - .offset: 24
1351 ; CHECK-NEXT: .size: 8
1352 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1353 ; CHECK-NEXT: - .offset: 32
1354 ; CHECK-NEXT: .size: 8
1355 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1356 ; CHECK-NEXT: - .offset: 40
1357 ; CHECK-NEXT: .size: 8
1358 ; CHECK-NEXT: .value_kind: hidden_none
1359 ; CHECK-NEXT: - .offset: 48
1360 ; CHECK-NEXT: .size: 8
1361 ; CHECK-NEXT: .value_kind: hidden_none
1362 ; CHECK-NEXT: - .offset: 56
1363 ; CHECK-NEXT: .size: 8
1364 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1365 ; CHECK: .language: OpenCL C
1366 ; CHECK-NEXT: .language_version:
1367 ; CHECK-NEXT: - 2
1368 ; CHECK-NEXT: - 0
1369 ; CHECK: .name: test_arg_struct_contains_ptr
1370 ; CHECK: .symbol: test_arg_struct_contains_ptr.kd
1371 define amdgpu_kernel void @test_arg_struct_contains_ptr(%struct.B %a) #0
1372 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !82
1373 !kernel_arg_base_type !82 !kernel_arg_type_qual !4 {
1374 ret void
1375 }
1377 ; CHECK: - .args:
1378 ; CHECK-NEXT: - .name: a
1379 ; CHECK-NEXT: .offset: 0
1380 ; CHECK-NEXT: .size: 16
1381 ; CHECK-NEXT: .type_name: 'global int addrspace(5)* __attribute__((ext_vector_type(2)))'
1382 ; CHECK-NEXT: .value_kind: by_value
1383 ; CHECK-NEXT: - .offset: 16
1384 ; CHECK-NEXT: .size: 8
1385 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1386 ; CHECK-NEXT: - .offset: 24
1387 ; CHECK-NEXT: .size: 8
1388 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1389 ; CHECK-NEXT: - .offset: 32
1390 ; CHECK-NEXT: .size: 8
1391 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1392 ; CHECK-NEXT: - .offset: 40
1393 ; CHECK-NEXT: .size: 8
1394 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1395 ; CHECK-NEXT: - .offset: 48
1396 ; CHECK-NEXT: .size: 8
1397 ; CHECK-NEXT: .value_kind: hidden_none
1398 ; CHECK-NEXT: - .offset: 56
1399 ; CHECK-NEXT: .size: 8
1400 ; CHECK-NEXT: .value_kind: hidden_none
1401 ; CHECK-NEXT: - .offset: 64
1402 ; CHECK-NEXT: .size: 8
1403 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1404 ; CHECK: .language: OpenCL C
1405 ; CHECK-NEXT: .language_version:
1406 ; CHECK-NEXT: - 2
1407 ; CHECK-NEXT: - 0
1408 ; CHECK: .name: test_arg_vector_of_ptr
1409 ; CHECK: .symbol: test_arg_vector_of_ptr.kd
1410 define amdgpu_kernel void @test_arg_vector_of_ptr(<2 x ptr addrspace(1)> %a) #0
1411 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !83
1412 !kernel_arg_base_type !83 !kernel_arg_type_qual !4 {
1413 ret void
1414 }
1416 ; CHECK: - .args:
1417 ; CHECK-NEXT: - .address_space: global
1418 ; CHECK-NEXT: .name: a
1419 ; CHECK-NEXT: .offset: 0
1420 ; CHECK-NEXT: .size: 8
1421 ; CHECK-NEXT: .type_name: clk_event_t
1422 ; CHECK-NEXT: .value_kind: global_buffer
1423 ; CHECK-NEXT: - .offset: 8
1424 ; CHECK-NEXT: .size: 8
1425 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1426 ; CHECK-NEXT: - .offset: 16
1427 ; CHECK-NEXT: .size: 8
1428 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1429 ; CHECK-NEXT: - .offset: 24
1430 ; CHECK-NEXT: .size: 8
1431 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1432 ; CHECK-NEXT: - .offset: 32
1433 ; CHECK-NEXT: .size: 8
1434 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1435 ; CHECK-NEXT: - .offset: 40
1436 ; CHECK-NEXT: .size: 8
1437 ; CHECK-NEXT: .value_kind: hidden_none
1438 ; CHECK-NEXT: - .offset: 48
1439 ; CHECK-NEXT: .size: 8
1440 ; CHECK-NEXT: .value_kind: hidden_none
1441 ; CHECK-NEXT: - .offset: 56
1442 ; CHECK-NEXT: .size: 8
1443 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1444 ; CHECK: .language: OpenCL C
1445 ; CHECK-NEXT: .language_version:
1446 ; CHECK-NEXT: - 2
1447 ; CHECK-NEXT: - 0
1448 ; CHECK: .name: test_arg_unknown_builtin_type
1449 ; CHECK: .symbol: test_arg_unknown_builtin_type.kd
1450 define amdgpu_kernel void @test_arg_unknown_builtin_type(
1451 ptr addrspace(1) %a) #0
1452 !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !84
1453 !kernel_arg_base_type !84 !kernel_arg_type_qual !4 {
1454 ret void
1455 }
1457 ; CHECK: - .args:
1458 ; CHECK-NEXT: - .address_space: global
1459 ; CHECK-NEXT: .name: a
1460 ; CHECK-NEXT: .offset: 0
1461 ; CHECK-NEXT: .size: 8
1462 ; CHECK-NEXT: .type_name: 'long addrspace(5)*'
1463 ; CHECK-NEXT: .value_kind: global_buffer
1464 ; CHECK-NEXT: - .address_space: local
1465 ; CHECK-NEXT: .name: b
1466 ; CHECK-NEXT: .offset: 8
1467 ; CHECK-NEXT: .pointee_align: 1
1468 ; CHECK-NEXT: .size: 4
1469 ; CHECK-NEXT: .type_name: 'char addrspace(5)*'
1470 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1471 ; CHECK-NEXT: - .address_space: local
1472 ; CHECK-NEXT: .name: c
1473 ; CHECK-NEXT: .offset: 12
1474 ; CHECK-NEXT: .pointee_align: 2
1475 ; CHECK-NEXT: .size: 4
1476 ; CHECK-NEXT: .type_name: 'char2 addrspace(5)*'
1477 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1478 ; CHECK-NEXT: - .address_space: local
1479 ; CHECK-NEXT: .name: d
1480 ; CHECK-NEXT: .offset: 16
1481 ; CHECK-NEXT: .pointee_align: 4
1482 ; CHECK-NEXT: .size: 4
1483 ; CHECK-NEXT: .type_name: 'char3 addrspace(5)*'
1484 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1485 ; CHECK-NEXT: - .address_space: local
1486 ; CHECK-NEXT: .name: e
1487 ; CHECK-NEXT: .offset: 20
1488 ; CHECK-NEXT: .pointee_align: 4
1489 ; CHECK-NEXT: .size: 4
1490 ; CHECK-NEXT: .type_name: 'char4 addrspace(5)*'
1491 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1492 ; CHECK-NEXT: - .address_space: local
1493 ; CHECK-NEXT: .name: f
1494 ; CHECK-NEXT: .offset: 24
1495 ; CHECK-NEXT: .pointee_align: 8
1496 ; CHECK-NEXT: .size: 4
1497 ; CHECK-NEXT: .type_name: 'char8 addrspace(5)*'
1498 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1499 ; CHECK-NEXT: - .address_space: local
1500 ; CHECK-NEXT: .name: g
1501 ; CHECK-NEXT: .offset: 28
1502 ; CHECK-NEXT: .pointee_align: 16
1503 ; CHECK-NEXT: .size: 4
1504 ; CHECK-NEXT: .type_name: 'char16 addrspace(5)*'
1505 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1506 ; CHECK-NEXT: - .address_space: local
1507 ; CHECK-NEXT: .name: h
1508 ; CHECK-NEXT: .offset: 32
1509 ; CHECK-NEXT: .pointee_align: 1
1510 ; CHECK-NEXT: .size: 4
1511 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1512 ; CHECK-NEXT: - .offset: 40
1513 ; CHECK-NEXT: .size: 8
1514 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1515 ; CHECK-NEXT: - .offset: 48
1516 ; CHECK-NEXT: .size: 8
1517 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1518 ; CHECK-NEXT: - .offset: 56
1519 ; CHECK-NEXT: .size: 8
1520 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1521 ; CHECK-NEXT: - .offset: 64
1522 ; CHECK-NEXT: .size: 8
1523 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1524 ; CHECK-NEXT: - .offset: 72
1525 ; CHECK-NEXT: .size: 8
1526 ; CHECK-NEXT: .value_kind: hidden_none
1527 ; CHECK-NEXT: - .offset: 80
1528 ; CHECK-NEXT: .size: 8
1529 ; CHECK-NEXT: .value_kind: hidden_none
1530 ; CHECK-NEXT: - .offset: 88
1531 ; CHECK-NEXT: .size: 8
1532 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1533 ; CHECK: .language: OpenCL C
1534 ; CHECK-NEXT: .language_version:
1535 ; CHECK-NEXT: - 2
1536 ; CHECK-NEXT: - 0
1537 ; CHECK: .name: test_pointee_align
1538 ; CHECK: .symbol: test_pointee_align.kd
1539 define amdgpu_kernel void @test_pointee_align(ptr addrspace(1) %a,
1540 ptr addrspace(3) %b,
1541 ptr addrspace(3) align 2 %c,
1542 ptr addrspace(3) align 4 %d,
1543 ptr addrspace(3) align 4 %e,
1544 ptr addrspace(3) align 8 %f,
1545 ptr addrspace(3) align 16 %g,
1546 ptr addrspace(3) %h) #0
1547 !kernel_arg_addr_space !91 !kernel_arg_access_qual !92 !kernel_arg_type !93
1548 !kernel_arg_base_type !93 !kernel_arg_type_qual !94 {
1549 ret void
1550 }
1552 ; CHECK: - .args:
1553 ; CHECK-NEXT: - .address_space: global
1554 ; CHECK-NEXT: .name: a
1555 ; CHECK-NEXT: .offset: 0
1556 ; CHECK-NEXT: .size: 8
1557 ; CHECK-NEXT: .type_name: 'long addrspace(5)*'
1558 ; CHECK-NEXT: .value_kind: global_buffer
1559 ; CHECK-NEXT: - .address_space: local
1560 ; CHECK-NEXT: .name: b
1561 ; CHECK-NEXT: .offset: 8
1562 ; CHECK-NEXT: .pointee_align: 8
1563 ; CHECK-NEXT: .size: 4
1564 ; CHECK-NEXT: .type_name: 'char addrspace(5)*'
1565 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1566 ; CHECK-NEXT: - .address_space: local
1567 ; CHECK-NEXT: .name: c
1568 ; CHECK-NEXT: .offset: 12
1569 ; CHECK-NEXT: .pointee_align: 32
1570 ; CHECK-NEXT: .size: 4
1571 ; CHECK-NEXT: .type_name: 'char2 addrspace(5)*'
1572 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1573 ; CHECK-NEXT: - .address_space: local
1574 ; CHECK-NEXT: .name: d
1575 ; CHECK-NEXT: .offset: 16
1576 ; CHECK-NEXT: .pointee_align: 64
1577 ; CHECK-NEXT: .size: 4
1578 ; CHECK-NEXT: .type_name: 'char3 addrspace(5)*'
1579 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1580 ; CHECK-NEXT: - .address_space: local
1581 ; CHECK-NEXT: .name: e
1582 ; CHECK-NEXT: .offset: 20
1583 ; CHECK-NEXT: .pointee_align: 256
1584 ; CHECK-NEXT: .size: 4
1585 ; CHECK-NEXT: .type_name: 'char4 addrspace(5)*'
1586 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1587 ; CHECK-NEXT: - .address_space: local
1588 ; CHECK-NEXT: .name: f
1589 ; CHECK-NEXT: .offset: 24
1590 ; CHECK-NEXT: .pointee_align: 128
1591 ; CHECK-NEXT: .size: 4
1592 ; CHECK-NEXT: .type_name: 'char8 addrspace(5)*'
1593 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1594 ; CHECK-NEXT: - .address_space: local
1595 ; CHECK-NEXT: .name: g
1596 ; CHECK-NEXT: .offset: 28
1597 ; CHECK-NEXT: .pointee_align: 1024
1598 ; CHECK-NEXT: .size: 4
1599 ; CHECK-NEXT: .type_name: 'char16 addrspace(5)*'
1600 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1601 ; CHECK-NEXT: - .address_space: local
1602 ; CHECK-NEXT: .name: h
1603 ; CHECK-NEXT: .offset: 32
1604 ; CHECK-NEXT: .pointee_align: 16
1605 ; CHECK-NEXT: .size: 4
1606 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1607 ; CHECK-NEXT: - .offset: 40
1608 ; CHECK-NEXT: .size: 8
1609 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1610 ; CHECK-NEXT: - .offset: 48
1611 ; CHECK-NEXT: .size: 8
1612 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1613 ; CHECK-NEXT: - .offset: 56
1614 ; CHECK-NEXT: .size: 8
1615 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1616 ; CHECK-NEXT: - .offset: 64
1617 ; CHECK-NEXT: .size: 8
1618 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1619 ; CHECK-NEXT: - .offset: 72
1620 ; CHECK-NEXT: .size: 8
1621 ; CHECK-NEXT: .value_kind: hidden_none
1622 ; CHECK-NEXT: - .offset: 80
1623 ; CHECK-NEXT: .size: 8
1624 ; CHECK-NEXT: .value_kind: hidden_none
1625 ; CHECK-NEXT: - .offset: 88
1626 ; CHECK-NEXT: .size: 8
1627 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1628 ; CHECK: .language: OpenCL C
1629 ; CHECK-NEXT: .language_version:
1630 ; CHECK-NEXT: - 2
1631 ; CHECK-NEXT: - 0
1632 ; CHECK: .name: test_pointee_align_attribute
1633 ; CHECK: .symbol: test_pointee_align_attribute.kd
1634 define amdgpu_kernel void @test_pointee_align_attribute(ptr addrspace(1) align 16 %a,
1635 ptr addrspace(3) align 8 %b,
1636 ptr addrspace(3) align 32 %c,
1637 ptr addrspace(3) align 64 %d,
1638 ptr addrspace(3) align 256 %e,
1639 ptr addrspace(3) align 128 %f,
1640 ptr addrspace(3) align 1024 %g,
1641 ptr addrspace(3) align 16 %h) #0
1642 !kernel_arg_addr_space !91 !kernel_arg_access_qual !92 !kernel_arg_type !93
1643 !kernel_arg_base_type !93 !kernel_arg_type_qual !94 {
1644 ret void
1645 }
1646 ; CHECK: - .args:
1647 ; CHECK-NEXT: - .name: arg
1648 ; CHECK-NEXT: .offset: 0
1649 ; CHECK-NEXT: .size: 25
1650 ; CHECK-NEXT: .type_name: __block_literal
1651 ; CHECK-NEXT: .value_kind: by_value
1652 ; CHECK-NEXT: - .offset: 32
1653 ; CHECK-NEXT: .size: 8
1654 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1655 ; CHECK-NEXT: - .offset: 40
1656 ; CHECK-NEXT: .size: 8
1657 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1658 ; CHECK-NEXT: - .offset: 48
1659 ; CHECK-NEXT: .size: 8
1660 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1661 ; CHECK-NEXT: - .offset: 56
1662 ; CHECK-NEXT: .size: 8
1663 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1664 ; CHECK-NEXT: - .offset: 64
1665 ; CHECK-NEXT: .size: 8
1666 ; CHECK-NEXT: .value_kind: hidden_none
1667 ; CHECK-NEXT: - .offset: 72
1668 ; CHECK-NEXT: .size: 8
1669 ; CHECK-NEXT: .value_kind: hidden_none
1670 ; CHECK-NEXT: - .offset: 80
1671 ; CHECK-NEXT: .size: 8
1672 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1673 ; CHECK: .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle
1674 ; CHECK: .language: OpenCL C
1675 ; CHECK-NEXT: .language_version:
1676 ; CHECK-NEXT: - 2
1677 ; CHECK-NEXT: - 0
1678 ; CHECK: .name: __test_block_invoke_kernel
1679 ; CHECK: .symbol: __test_block_invoke_kernel.kd
1680 define amdgpu_kernel void @__test_block_invoke_kernel(
1681 <{ i32, i32, ptr, ptr addrspace(1), i8 }> %arg) #1
1682 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !110
1683 !kernel_arg_base_type !110 !kernel_arg_type_qual !4 {
1684 ret void
1685 }
1687 ; CHECK: - .args:
1688 ; CHECK-NEXT: - .name: a
1689 ; CHECK-NEXT: .offset: 0
1690 ; CHECK-NEXT: .size: 1
1691 ; CHECK-NEXT: .type_name: char
1692 ; CHECK-NEXT: .value_kind: by_value
1693 ; CHECK-NEXT: - .offset: 8
1694 ; CHECK-NEXT: .size: 8
1695 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1696 ; CHECK-NEXT: - .offset: 16
1697 ; CHECK-NEXT: .size: 8
1698 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1699 ; CHECK-NEXT: - .offset: 24
1700 ; CHECK-NEXT: .size: 8
1701 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1702 ; CHECK-NEXT: - .offset: 32
1703 ; CHECK-NEXT: .size: 8
1704 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1705 ; CHECK-NEXT: - .offset: 40
1706 ; CHECK-NEXT: .size: 8
1707 ; CHECK-NEXT: .value_kind: hidden_default_queue
1708 ; CHECK-NEXT: - .offset: 48
1709 ; CHECK-NEXT: .size: 8
1710 ; CHECK-NEXT: .value_kind: hidden_completion_action
1711 ; CHECK-NEXT: - .offset: 56
1712 ; CHECK-NEXT: .size: 8
1713 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1714 ; CHECK: .language: OpenCL C
1715 ; CHECK-NEXT: .language_version:
1716 ; CHECK-NEXT: - 2
1717 ; CHECK-NEXT: - 0
1718 ; CHECK: .name: test_enqueue_kernel_caller
1719 ; CHECK: .symbol: test_enqueue_kernel_caller.kd
1720 define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #2
1721 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
1722 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
1723 ret void
1724 }
1726 ; CHECK: - .args:
1727 ; CHECK-NEXT: - .name: ptr
1728 ; CHECK-NEXT: .offset: 0
1729 ; CHECK-NEXT: .size: 8
1730 ; CHECK-NEXT: .value_kind: global_buffer
1731 ; CHECK: .name: unknown_addrspace_kernarg
1732 ; CHECK: .symbol: unknown_addrspace_kernarg.kd
1733 define amdgpu_kernel void @unknown_addrspace_kernarg(ptr addrspace(12345) %ptr) #0 {
1734 ret void
1735 }
1737 ; CHECK: amdhsa.printf:
1738 ; CHECK-NEXT: - '1:1:4:%d\n'
1739 ; CHECK-NEXT: - '2:1:8:%g\n'
1740 ; CHECK: amdhsa.version:
1741 ; CHECK-NEXT: - 1
1742 ; CHECK-NEXT: - 0
1744 attributes #0 = { optnone noinline "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-implicitarg-num-bytes"="56" }
1745 attributes #1 = { optnone noinline "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-implicitarg-num-bytes"="56" "runtime-handle"="__test_block_invoke_kernel_runtime_handle" }
1746 attributes #2 = { optnone noinline "amdgpu-implicitarg-num-bytes"="56" }
1748 !llvm.module.flags = !{!0}
1749 !0 = !{i32 1, !"amdgpu_code_object_version", i32 300}
1751 !llvm.printf.fmts = !{!100, !101}
1753 !1 = !{i32 0}
1754 !2 = !{!"none"}
1755 !3 = !{!"int"}
1756 !4 = !{!""}
1757 !5 = !{i32 undef, i32 1}
1758 !6 = !{i32 1, i32 2, i32 4}
1759 !7 = !{<4 x i32> undef, i32 0}
1760 !8 = !{i32 8, i32 16, i32 32}
1761 !9 = !{!"char"}
1762 !10 = !{!"ushort2"}
1763 !11 = !{!"int3"}
1764 !12 = !{!"ulong4"}
1765 !13 = !{!"half8"}
1766 !14 = !{!"float16"}
1767 !15 = !{!"double16"}
1768 !16 = !{!"int addrspace(5)*"}
1769 !17 = !{!"image2d_t"}
1770 !18 = !{!"sampler_t"}
1771 !19 = !{!"queue_t"}
1772 !20 = !{!"struct A"}
1773 !21 = !{!"i128"}
1774 !22 = !{i32 0, i32 0, i32 0}
1775 !23 = !{!"none", !"none", !"none"}
1776 !24 = !{!"int", !"short2", !"char3"}
1777 !25 = !{!"", !"", !""}
1778 !26 = !{half undef, i32 1}
1779 !27 = !{float undef, i32 1}
1780 !28 = !{double undef, i32 1}
1781 !29 = !{i8 undef, i32 1}
1782 !30 = !{i16 undef, i32 1}
1783 !31 = !{i64 undef, i32 1}
1784 !32 = !{ptr addrspace(5) undef, i32 1}
1785 !50 = !{i32 1, i32 2, i32 3}
1786 !51 = !{!"int addrspace(5)*", !"int addrspace(5)*", !"int addrspace(5)*"}
1787 !60 = !{i32 1, i32 1, i32 1}
1788 !61 = !{!"read_only", !"write_only", !"read_write"}
1789 !62 = !{!"image1d_t", !"image2d_t", !"image3d_t"}
1790 !70 = !{!"volatile", !"const restrict", !"pipe"}
1791 !80 = !{!"int addrspace(5)* addrspace(5)*"}
1792 !81 = !{i32 1}
1793 !82 = !{!"struct B"}
1794 !83 = !{!"global int addrspace(5)* __attribute__((ext_vector_type(2)))"}
1795 !84 = !{!"clk_event_t"}
1796 !opencl.ocl.version = !{!90}
1797 !90 = !{i32 2, i32 0}
1798 !91 = !{i32 0, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3}
1799 !92 = !{!"none", !"none", !"none", !"none", !"none", !"none", !"none"}
1800 !93 = !{!"long addrspace(5)*", !"char addrspace(5)*", !"char2 addrspace(5)*", !"char3 addrspace(5)*", !"char4 addrspace(5)*", !"char8 addrspace(5)*", !"char16 addrspace(5)*"}
1801 !94 = !{!"", !"", !"", !"", !"", !"", !""}
1802 !100 = !{!"1:1:4:%d\5Cn"}
1803 !101 = !{!"2:1:8:%g\5Cn"}
1804 !110 = !{!"__block_literal"}
1805 !111 = !{!"char", !"char"}
1807 ; PARSER: AMDGPU HSA Metadata Parser Test: PASS