Lines Matching refs:CHECK
1 …iletype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK --check-prefix=GFX…
2 …iletype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK --check-prefix=GFX…
3 …iletype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK --check-prefix=GFX…
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: - .address_space: global
37 ; CHECK-NEXT: .offset: 32
38 ; CHECK-NEXT: .size: 8
39 ; CHECK-NOT: .value_kind: hidden_default_queue
40 ; CHECK-NOT: .value_kind: hidden_completion_action
41 ; CHECK-NOT: .value_kind: hidden_hostcall_buffer
42 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
43 ; CHECK: .value_kind: hidden_multigrid_sync_arg
44 ; CHECK: .language: OpenCL C
45 ; CHECK-NEXT: .language_version:
46 ; CHECK-NEXT: - 2
47 ; CHECK-NEXT: - 0
48 ; CHECK: .name: test_char
49 ; CHECK: .symbol: test_char.kd
56 ; CHECK: - .args:
57 ; CHECK-NEXT: - .name: a
58 ; CHECK-NEXT: .offset: 0
59 ; CHECK-NEXT: .size: 1
60 ; CHECK-NEXT: .type_name: char
61 ; CHECK-NEXT: .value_kind: by_value
62 ; CHECK-NEXT: - .offset: 8
63 ; CHECK-NEXT: .size: 8
64 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
65 ; CHECK-NEXT: - .offset: 16
66 ; CHECK-NEXT: .size: 8
67 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
68 ; CHECK-NEXT: - .offset: 24
69 ; CHECK-NEXT: .size: 8
70 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
71 ; CHECK-NEXT: - .address_space: global
72 ; CHECK-NEXT: .offset: 32
73 ; CHECK-NEXT: .size: 8
74 ; CHECK-NOT: .value_kind: hidden_default_queue
75 ; CHECK-NOT: .value_kind: hidden_completion_action
76 ; CHECK-NOT: .value_kind: hidden_hostcall_buffer
77 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
78 ; CHECK: .value_kind: hidden_multigrid_sync_arg
79 ; CHECK: .language: OpenCL C
80 ; CHECK-NEXT: .language_version:
81 ; CHECK-NEXT: - 2
82 ; CHECK-NEXT: - 0
83 ; CHECK: .name: test_char_byref_constant
84 ; CHECK: .symbol: test_char_byref_constant.kd
91 ; CHECK: - .args:
92 ; CHECK-NEXT: - .offset: 0
93 ; CHECK-NEXT: .size: 1
94 ; CHECK-NEXT: .type_name: char
95 ; CHECK-NEXT: .value_kind: by_value
96 ; CHECK-NEXT: - .name: a
97 ; CHECK-NEXT: .offset: 512
98 ; CHECK-NEXT: .size: 1
99 ; CHECK-NEXT: .type_name: char
100 ; CHECK-NEXT: .value_kind: by_value
101 ; CHECK-NEXT: - .offset: 520
102 ; CHECK-NEXT: .size: 8
103 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
104 ; CHECK-NEXT: - .offset: 528
105 ; CHECK-NEXT: .size: 8
106 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
107 ; CHECK-NEXT: - .offset: 536
108 ; CHECK-NEXT: .size: 8
109 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
110 ; CHECK-NEXT: - .address_space: global
111 ; CHECK-NEXT: .offset: 544
112 ; CHECK-NEXT: .size: 8
113 ; CHECK-NOT: .value_kind: hidden_default_queue
114 ; CHECK-NOT: .value_kind: hidden_completion_action
115 ; CHECK-NOT: .value_kind: hidden_hostcall_buffer
116 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
117 ; CHECK: .value_kind: hidden_multigrid_sync_arg
118 ; CHECK: .language: OpenCL C
119 ; CHECK-NEXT: .language_version:
120 ; CHECK-NEXT: - 2
121 ; CHECK-NEXT: - 0
122 ; CHECK: .name: test_char_byref_constant_align512
123 ; CHECK: .symbol: test_char_byref_constant_align512.kd
130 ; CHECK: - .args:
131 ; CHECK-NEXT: - .name: a
132 ; CHECK-NEXT: .offset: 0
133 ; CHECK-NEXT: .size: 4
134 ; CHECK-NEXT: .type_name: ushort2
135 ; CHECK-NEXT: .value_kind: by_value
136 ; CHECK-NEXT: - .offset: 8
137 ; CHECK-NEXT: .size: 8
138 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
139 ; CHECK-NEXT: - .offset: 16
140 ; CHECK-NEXT: .size: 8
141 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
142 ; CHECK-NEXT: - .offset: 24
143 ; CHECK-NEXT: .size: 8
144 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
145 ; CHECK-NEXT: - .address_space: global
146 ; CHECK-NEXT: .offset: 32
147 ; CHECK-NEXT: .size: 8
148 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
149 ; CHECK-NEXT: - .address_space: global
150 ; CHECK-NEXT: .offset: 40
151 ; CHECK-NEXT: .size: 8
152 ; CHECK-NEXT: .value_kind: hidden_none
153 ; CHECK-NEXT: - .address_space: global
154 ; CHECK-NEXT: .offset: 48
155 ; CHECK-NEXT: .size: 8
156 ; CHECK-NEXT: .value_kind: hidden_none
157 ; CHECK-NEXT: - .address_space: global
158 ; CHECK-NEXT: .offset: 56
159 ; CHECK-NEXT: .size: 8
160 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
161 ; CHECK: .language: OpenCL C
162 ; CHECK-NEXT: .language_version:
163 ; CHECK-NEXT: - 2
164 ; CHECK-NEXT: - 0
165 ; CHECK: .name: test_ushort2
166 ; CHECK: .symbol: test_ushort2.kd
173 ; CHECK: - .args:
174 ; CHECK-NEXT: - .name: a
175 ; CHECK-NEXT: .offset: 0
176 ; CHECK-NEXT: .size: 16
177 ; CHECK-NEXT: .type_name: int3
178 ; CHECK-NEXT: .value_kind: by_value
179 ; CHECK-NEXT: - .offset: 16
180 ; CHECK-NEXT: .size: 8
181 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
182 ; CHECK-NEXT: - .offset: 24
183 ; CHECK-NEXT: .size: 8
184 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
185 ; CHECK-NEXT: - .offset: 32
186 ; CHECK-NEXT: .size: 8
187 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
188 ; CHECK-NEXT: - .address_space: global
189 ; CHECK-NEXT: .offset: 40
190 ; CHECK-NEXT: .size: 8
191 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
192 ; CHECK-NEXT: - .address_space: global
193 ; CHECK-NEXT: .offset: 48
194 ; CHECK-NEXT: .size: 8
195 ; CHECK-NEXT: .value_kind: hidden_none
196 ; CHECK-NEXT: - .address_space: global
197 ; CHECK-NEXT: .offset: 56
198 ; CHECK-NEXT: .size: 8
199 ; CHECK-NEXT: .value_kind: hidden_none
200 ; CHECK-NEXT: - .address_space: global
201 ; CHECK-NEXT: .offset: 64
202 ; CHECK-NEXT: .size: 8
203 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
204 ; CHECK: .language: OpenCL C
205 ; CHECK-NEXT: .language_version:
206 ; CHECK-NEXT: - 2
207 ; CHECK-NEXT: - 0
208 ; CHECK: .name: test_int3
209 ; CHECK: .symbol: test_int3.kd
216 ; CHECK: - .args:
217 ; CHECK-NEXT: - .name: a
218 ; CHECK-NEXT: .offset: 0
219 ; CHECK-NEXT: .size: 32
220 ; CHECK-NEXT: .type_name: ulong4
221 ; CHECK-NEXT: .value_kind: by_value
222 ; CHECK-NEXT: - .offset: 32
223 ; CHECK-NEXT: .size: 8
224 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
225 ; CHECK-NEXT: - .offset: 40
226 ; CHECK-NEXT: .size: 8
227 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
228 ; CHECK-NEXT: - .offset: 48
229 ; CHECK-NEXT: .size: 8
230 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
231 ; CHECK-NEXT: - .address_space: global
232 ; CHECK-NEXT: .offset: 56
233 ; CHECK-NEXT: .size: 8
234 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
235 ; CHECK-NEXT: - .address_space: global
236 ; CHECK-NEXT: .offset: 64
237 ; CHECK-NEXT: .size: 8
238 ; CHECK-NEXT: .value_kind: hidden_none
239 ; CHECK-NEXT: - .address_space: global
240 ; CHECK-NEXT: .offset: 72
241 ; CHECK-NEXT: .size: 8
242 ; CHECK-NEXT: .value_kind: hidden_none
243 ; CHECK-NEXT: - .address_space: global
244 ; CHECK-NEXT: .offset: 80
245 ; CHECK-NEXT: .size: 8
246 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
247 ; CHECK: .language: OpenCL C
248 ; CHECK-NEXT: .language_version:
249 ; CHECK-NEXT: - 2
250 ; CHECK-NEXT: - 0
251 ; CHECK: .name: test_ulong4
252 ; CHECK: .symbol: test_ulong4.kd
259 ; CHECK: - .args:
260 ; CHECK-NEXT: - .name: a
261 ; CHECK-NEXT: .offset: 0
262 ; CHECK-NEXT: .size: 16
263 ; CHECK-NEXT: .type_name: half8
264 ; CHECK-NEXT: .value_kind: by_value
265 ; CHECK-NEXT: - .offset: 16
266 ; CHECK-NEXT: .size: 8
267 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
268 ; CHECK-NEXT: - .offset: 24
269 ; CHECK-NEXT: .size: 8
270 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
271 ; CHECK-NEXT: - .offset: 32
272 ; CHECK-NEXT: .size: 8
273 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
274 ; CHECK-NEXT: - .address_space: global
275 ; CHECK-NEXT: .offset: 40
276 ; CHECK-NEXT: .size: 8
277 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
278 ; CHECK-NEXT: - .address_space: global
279 ; CHECK-NEXT: .offset: 48
280 ; CHECK-NEXT: .size: 8
281 ; CHECK-NEXT: .value_kind: hidden_none
282 ; CHECK-NEXT: - .address_space: global
283 ; CHECK-NEXT: .offset: 56
284 ; CHECK-NEXT: .size: 8
285 ; CHECK-NEXT: .value_kind: hidden_none
286 ; CHECK-NEXT: - .address_space: global
287 ; CHECK-NEXT: .offset: 64
288 ; CHECK-NEXT: .size: 8
289 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
290 ; CHECK: .language: OpenCL C
291 ; CHECK-NEXT: .language_version:
292 ; CHECK-NEXT: - 2
293 ; CHECK-NEXT: - 0
294 ; CHECK: .name: test_half8
295 ; CHECK: .symbol: test_half8.kd
302 ; CHECK: - .args:
303 ; CHECK-NEXT: - .name: a
304 ; CHECK-NEXT: .offset: 0
305 ; CHECK-NEXT: .size: 64
306 ; CHECK-NEXT: .type_name: float16
307 ; CHECK-NEXT: .value_kind: by_value
308 ; CHECK-NEXT: - .offset: 64
309 ; CHECK-NEXT: .size: 8
310 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
311 ; CHECK-NEXT: - .offset: 72
312 ; CHECK-NEXT: .size: 8
313 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
314 ; CHECK-NEXT: - .offset: 80
315 ; CHECK-NEXT: .size: 8
316 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
317 ; CHECK-NEXT: - .address_space: global
318 ; CHECK-NEXT: .offset: 88
319 ; CHECK-NEXT: .size: 8
320 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
321 ; CHECK-NEXT: - .address_space: global
322 ; CHECK-NEXT: .offset: 96
323 ; CHECK-NEXT: .size: 8
324 ; CHECK-NEXT: .value_kind: hidden_none
325 ; CHECK-NEXT: - .address_space: global
326 ; CHECK-NEXT: .offset: 104
327 ; CHECK-NEXT: .size: 8
328 ; CHECK-NEXT: .value_kind: hidden_none
329 ; CHECK-NEXT: - .address_space: global
330 ; CHECK-NEXT: .offset: 112
331 ; CHECK-NEXT: .size: 8
332 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
333 ; CHECK: .language: OpenCL C
334 ; CHECK-NEXT: .language_version:
335 ; CHECK-NEXT: - 2
336 ; CHECK-NEXT: - 0
337 ; CHECK: .name: test_float16
338 ; CHECK: .symbol: test_float16.kd
345 ; CHECK: - .args:
346 ; CHECK-NEXT: - .name: a
347 ; CHECK-NEXT: .offset: 0
348 ; CHECK-NEXT: .size: 128
349 ; CHECK-NEXT: .type_name: double16
350 ; CHECK-NEXT: .value_kind: by_value
351 ; CHECK-NEXT: - .offset: 128
352 ; CHECK-NEXT: .size: 8
353 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
354 ; CHECK-NEXT: - .offset: 136
355 ; CHECK-NEXT: .size: 8
356 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
357 ; CHECK-NEXT: - .offset: 144
358 ; CHECK-NEXT: .size: 8
359 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
360 ; CHECK-NEXT: - .address_space: global
361 ; CHECK-NEXT: .offset: 152
362 ; CHECK-NEXT: .size: 8
363 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
364 ; CHECK-NEXT: - .address_space: global
365 ; CHECK-NEXT: .offset: 160
366 ; CHECK-NEXT: .size: 8
367 ; CHECK-NEXT: .value_kind: hidden_none
368 ; CHECK-NEXT: - .address_space: global
369 ; CHECK-NEXT: .offset: 168
370 ; CHECK-NEXT: .size: 8
371 ; CHECK-NEXT: .value_kind: hidden_none
372 ; CHECK-NEXT: - .address_space: global
373 ; CHECK-NEXT: .offset: 176
374 ; CHECK-NEXT: .size: 8
375 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
376 ; CHECK: .language: OpenCL C
377 ; CHECK-NEXT: .language_version:
378 ; CHECK-NEXT: - 2
379 ; CHECK-NEXT: - 0
380 ; CHECK: .name: test_double16
381 ; CHECK: .symbol: test_double16.kd
388 ; CHECK: - .args:
389 ; CHECK-NEXT: - .address_space: global
390 ; CHECK-NEXT: .name: a
391 ; CHECK-NEXT: .offset: 0
392 ; CHECK-NEXT: .size: 8
393 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
394 ; CHECK-NEXT: .value_kind: global_buffer
395 ; CHECK-NEXT: - .offset: 8
396 ; CHECK-NEXT: .size: 8
397 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
398 ; CHECK-NEXT: - .offset: 16
399 ; CHECK-NEXT: .size: 8
400 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
401 ; CHECK-NEXT: - .offset: 24
402 ; CHECK-NEXT: .size: 8
403 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
404 ; CHECK-NEXT: - .address_space: global
405 ; CHECK-NEXT: .offset: 32
406 ; CHECK-NEXT: .size: 8
407 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
408 ; CHECK-NEXT: - .address_space: global
409 ; CHECK-NEXT: .offset: 40
410 ; CHECK-NEXT: .size: 8
411 ; CHECK-NEXT: .value_kind: hidden_none
412 ; CHECK-NEXT: - .address_space: global
413 ; CHECK-NEXT: .offset: 48
414 ; CHECK-NEXT: .size: 8
415 ; CHECK-NEXT: .value_kind: hidden_none
416 ; CHECK-NEXT: - .address_space: global
417 ; CHECK-NEXT: .offset: 56
418 ; CHECK-NEXT: .size: 8
419 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
420 ; CHECK: .language: OpenCL C
421 ; CHECK-NEXT: .language_version:
422 ; CHECK-NEXT: - 2
423 ; CHECK-NEXT: - 0
424 ; CHECK: .name: test_pointer
425 ; CHECK: .symbol: test_pointer.kd
432 ; CHECK: - .args:
433 ; CHECK-NEXT: - .address_space: global
434 ; CHECK-NEXT: .name: a
435 ; CHECK-NEXT: .offset: 0
436 ; CHECK-NEXT: .size: 8
437 ; CHECK-NEXT: .type_name: image2d_t
438 ; CHECK-NEXT: .value_kind: image
439 ; CHECK-NEXT: - .offset: 8
440 ; CHECK-NEXT: .size: 8
441 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
442 ; CHECK-NEXT: - .offset: 16
443 ; CHECK-NEXT: .size: 8
444 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
445 ; CHECK-NEXT: - .offset: 24
446 ; CHECK-NEXT: .size: 8
447 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
448 ; CHECK-NEXT: - .address_space: global
449 ; CHECK-NEXT: .offset: 32
450 ; CHECK-NEXT: .size: 8
451 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
452 ; CHECK-NEXT: - .address_space: global
453 ; CHECK-NEXT: .offset: 40
454 ; CHECK-NEXT: .size: 8
455 ; CHECK-NEXT: .value_kind: hidden_none
456 ; CHECK-NEXT: - .address_space: global
457 ; CHECK-NEXT: .offset: 48
458 ; CHECK-NEXT: .size: 8
459 ; CHECK-NEXT: .value_kind: hidden_none
460 ; CHECK-NEXT: - .address_space: global
461 ; CHECK-NEXT: .offset: 56
462 ; CHECK-NEXT: .size: 8
463 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
464 ; CHECK: .language: OpenCL C
465 ; CHECK-NEXT: .language_version:
466 ; CHECK-NEXT: - 2
467 ; CHECK-NEXT: - 0
468 ; CHECK: .name: test_image
469 ; CHECK: .symbol: test_image.kd
476 ; CHECK: - .args:
477 ; CHECK-NEXT: - .name: a
478 ; CHECK-NEXT: .offset: 0
479 ; CHECK-NEXT: .size: 4
480 ; CHECK-NEXT: .type_name: sampler_t
481 ; CHECK-NEXT: .value_kind: sampler
482 ; CHECK-NEXT: - .offset: 8
483 ; CHECK-NEXT: .size: 8
484 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
485 ; CHECK-NEXT: - .offset: 16
486 ; CHECK-NEXT: .size: 8
487 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
488 ; CHECK-NEXT: - .offset: 24
489 ; CHECK-NEXT: .size: 8
490 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
491 ; CHECK-NEXT: - .address_space: global
492 ; CHECK-NEXT: .offset: 32
493 ; CHECK-NEXT: .size: 8
494 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
495 ; CHECK-NEXT: - .address_space: global
496 ; CHECK-NEXT: .offset: 40
497 ; CHECK-NEXT: .size: 8
498 ; CHECK-NEXT: .value_kind: hidden_none
499 ; CHECK-NEXT: - .address_space: global
500 ; CHECK-NEXT: .offset: 48
501 ; CHECK-NEXT: .size: 8
502 ; CHECK-NEXT: .value_kind: hidden_none
503 ; CHECK-NEXT: - .address_space: global
504 ; CHECK-NEXT: .offset: 56
505 ; CHECK-NEXT: .size: 8
506 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
507 ; CHECK: .language: OpenCL C
508 ; CHECK-NEXT: .language_version:
509 ; CHECK-NEXT: - 2
510 ; CHECK-NEXT: - 0
511 ; CHECK: .name: test_sampler
512 ; CHECK: .symbol: test_sampler.kd
519 ; CHECK: - .args:
520 ; CHECK-NEXT: - .address_space: global
521 ; CHECK-NEXT: .name: a
522 ; CHECK-NEXT: .offset: 0
523 ; CHECK-NEXT: .size: 8
524 ; CHECK-NEXT: .type_name: queue_t
525 ; CHECK-NEXT: .value_kind: queue
526 ; CHECK-NEXT: - .offset: 8
527 ; CHECK-NEXT: .size: 8
528 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
529 ; CHECK-NEXT: - .offset: 16
530 ; CHECK-NEXT: .size: 8
531 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
532 ; CHECK-NEXT: - .offset: 24
533 ; CHECK-NEXT: .size: 8
534 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
535 ; CHECK-NEXT: - .address_space: global
536 ; CHECK-NEXT: .offset: 32
537 ; CHECK-NEXT: .size: 8
538 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
539 ; CHECK-NEXT: - .address_space: global
540 ; CHECK-NEXT: .offset: 40
541 ; CHECK-NEXT: .size: 8
542 ; CHECK-NEXT: .value_kind: hidden_none
543 ; CHECK-NEXT: - .address_space: global
544 ; CHECK-NEXT: .offset: 48
545 ; CHECK-NEXT: .size: 8
546 ; CHECK-NEXT: .value_kind: hidden_none
547 ; CHECK-NEXT: - .address_space: global
548 ; CHECK-NEXT: .offset: 56
549 ; CHECK-NEXT: .size: 8
550 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
551 ; CHECK: .language: OpenCL C
552 ; CHECK-NEXT: .language_version:
553 ; CHECK-NEXT: - 2
554 ; CHECK-NEXT: - 0
555 ; CHECK: .name: test_queue
556 ; CHECK: .symbol: test_queue.kd
563 ; CHECK: - .args:
564 ; CHECK-NEXT: .name: a
565 ; CHECK-NEXT: .offset: 0
566 ; CHECK-NEXT: .size: 8
567 ; CHECK-NEXT: .type_name: struct A
568 ; CHECK-NEXT: .value_kind: by_value
569 ; CHECK-NEXT: - .offset: 8
570 ; CHECK-NEXT: .size: 8
571 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
572 ; CHECK-NEXT: - .offset: 16
573 ; CHECK-NEXT: .size: 8
574 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
575 ; CHECK-NEXT: - .offset: 24
576 ; CHECK-NEXT: .size: 8
577 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
578 ; CHECK-NEXT: - .address_space: global
579 ; CHECK-NEXT: .offset: 32
580 ; CHECK-NEXT: .size: 8
581 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
582 ; CHECK-NEXT: - .address_space: global
583 ; CHECK-NEXT: .offset: 40
584 ; CHECK-NEXT: .size: 8
585 ; CHECK-NEXT: .value_kind: hidden_none
586 ; CHECK-NEXT: - .address_space: global
587 ; CHECK-NEXT: .offset: 48
588 ; CHECK-NEXT: .size: 8
589 ; CHECK-NEXT: .value_kind: hidden_none
590 ; CHECK-NEXT: - .address_space: global
591 ; CHECK-NEXT: .offset: 56
592 ; CHECK-NEXT: .size: 8
593 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
594 ; CHECK: .language: OpenCL C
595 ; CHECK-NEXT: .language_version:
596 ; CHECK-NEXT: - 2
597 ; CHECK-NEXT: - 0
598 ; CHECK: .name: test_struct
599 ; CHECK: .symbol: test_struct.kd
606 ; CHECK: - .args:
607 ; CHECK-NEXT: .name: a
608 ; CHECK-NEXT: .offset: 0
609 ; CHECK-NEXT: .size: 8
610 ; CHECK-NEXT: .type_name: struct A
611 ; CHECK-NEXT: .value_kind: by_value
612 ; CHECK-NEXT: - .offset: 8
613 ; CHECK-NEXT: .size: 8
614 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
615 ; CHECK-NEXT: - .offset: 16
616 ; CHECK-NEXT: .size: 8
617 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
618 ; CHECK-NEXT: - .offset: 24
619 ; CHECK-NEXT: .size: 8
620 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
621 ; CHECK-NEXT: - .address_space: global
622 ; CHECK-NEXT: .offset: 32
623 ; CHECK-NEXT: .size: 8
624 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
625 ; CHECK-NEXT: - .address_space: global
626 ; CHECK-NEXT: .offset: 40
627 ; CHECK-NEXT: .size: 8
628 ; CHECK-NEXT: .value_kind: hidden_none
629 ; CHECK-NEXT: - .address_space: global
630 ; CHECK-NEXT: .offset: 48
631 ; CHECK-NEXT: .size: 8
632 ; CHECK-NEXT: .value_kind: hidden_none
633 ; CHECK-NEXT: - .address_space: global
634 ; CHECK-NEXT: .offset: 56
635 ; CHECK-NEXT: .size: 8
636 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
637 ; CHECK: .language: OpenCL C
638 ; CHECK-NEXT: .language_version:
639 ; CHECK-NEXT: - 2
640 ; CHECK-NEXT: - 0
641 ; CHECK: .name: test_struct_byref_constant
642 ; CHECK: .symbol: test_struct_byref_constant.kd
649 ; CHECK: - .args:
650 ; CHECK-NEXT: .name: a
651 ; CHECK-NEXT: .offset: 0
652 ; CHECK-NEXT: .size: 32
653 ; CHECK-NEXT: .type_name: struct A
654 ; CHECK-NEXT: .value_kind: by_value
655 ; CHECK-NEXT: - .offset: 32
656 ; CHECK-NEXT: .size: 8
657 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
658 ; CHECK-NEXT: - .offset: 40
659 ; CHECK-NEXT: .size: 8
660 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
661 ; CHECK-NEXT: - .offset: 48
662 ; CHECK-NEXT: .size: 8
663 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
664 ; CHECK-NEXT: - .address_space: global
665 ; CHECK-NEXT: .offset: 56
666 ; CHECK-NEXT: .size: 8
667 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
668 ; CHECK-NEXT: - .address_space: global
669 ; CHECK-NEXT: .offset: 64
670 ; CHECK-NEXT: .size: 8
671 ; CHECK-NEXT: .value_kind: hidden_none
672 ; CHECK-NEXT: - .address_space: global
673 ; CHECK-NEXT: .offset: 72
674 ; CHECK-NEXT: .size: 8
675 ; CHECK-NEXT: .value_kind: hidden_none
676 ; CHECK-NEXT: - .address_space: global
677 ; CHECK-NEXT: .offset: 80
678 ; CHECK-NEXT: .size: 8
679 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
680 ; CHECK: .language: OpenCL C
681 ; CHECK-NEXT: .language_version:
682 ; CHECK-NEXT: - 2
683 ; CHECK-NEXT: - 0
684 ; CHECK: .name: test_array
685 ; CHECK: .symbol: test_array.kd
692 ; CHECK: - .args:
693 ; CHECK-NEXT: .name: a
694 ; CHECK-NEXT: .offset: 0
695 ; CHECK-NEXT: .size: 32
696 ; CHECK-NEXT: .type_name: struct A
697 ; CHECK-NEXT: .value_kind: by_value
698 ; CHECK-NEXT: - .offset: 32
699 ; CHECK-NEXT: .size: 8
700 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
701 ; CHECK-NEXT: - .offset: 40
702 ; CHECK-NEXT: .size: 8
703 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
704 ; CHECK-NEXT: - .offset: 48
705 ; CHECK-NEXT: .size: 8
706 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
707 ; CHECK-NEXT: - .address_space: global
708 ; CHECK-NEXT: .offset: 56
709 ; CHECK-NEXT: .size: 8
710 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
711 ; CHECK-NEXT: - .address_space: global
712 ; CHECK-NEXT: .offset: 64
713 ; CHECK-NEXT: .size: 8
714 ; CHECK-NEXT: .value_kind: hidden_none
715 ; CHECK-NEXT: - .address_space: global
716 ; CHECK-NEXT: .offset: 72
717 ; CHECK-NEXT: .size: 8
718 ; CHECK-NEXT: .value_kind: hidden_none
719 ; CHECK-NEXT: - .address_space: global
720 ; CHECK-NEXT: .offset: 80
721 ; CHECK-NEXT: .size: 8
722 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
723 ; CHECK: .language: OpenCL C
724 ; CHECK-NEXT: .language_version:
725 ; CHECK-NEXT: - 2
726 ; CHECK-NEXT: - 0
727 ; CHECK: .name: test_array_byref_constant
728 ; CHECK: .symbol: test_array_byref_constant.kd
735 ; CHECK: - .args:
736 ; CHECK-NEXT: - .name: a
737 ; CHECK-NEXT: .offset: 0
738 ; CHECK-NEXT: .size: 16
739 ; CHECK-NEXT: .type_name: i128
740 ; CHECK-NEXT: .value_kind: by_value
741 ; CHECK-NEXT: - .offset: 16
742 ; CHECK-NEXT: .size: 8
743 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
744 ; CHECK-NEXT: - .offset: 24
745 ; CHECK-NEXT: .size: 8
746 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
747 ; CHECK-NEXT: - .offset: 32
748 ; CHECK-NEXT: .size: 8
749 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
750 ; CHECK-NEXT: - .address_space: global
751 ; CHECK-NEXT: .offset: 40
752 ; CHECK-NEXT: .size: 8
753 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
754 ; CHECK-NEXT: - .address_space: global
755 ; CHECK-NEXT: .offset: 48
756 ; CHECK-NEXT: .size: 8
757 ; CHECK-NEXT: .value_kind: hidden_none
758 ; CHECK-NEXT: - .address_space: global
759 ; CHECK-NEXT: .offset: 56
760 ; CHECK-NEXT: .size: 8
761 ; CHECK-NEXT: .value_kind: hidden_none
762 ; CHECK-NEXT: - .address_space: global
763 ; CHECK-NEXT: .offset: 64
764 ; CHECK-NEXT: .size: 8
765 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
766 ; CHECK: .language: OpenCL C
767 ; CHECK-NEXT: .language_version:
768 ; CHECK-NEXT: - 2
769 ; CHECK-NEXT: - 0
770 ; CHECK: .name: test_i128
771 ; CHECK: .symbol: test_i128.kd
778 ; CHECK: - .args:
779 ; CHECK-NEXT: - .name: a
780 ; CHECK-NEXT: .offset: 0
781 ; CHECK-NEXT: .size: 4
782 ; CHECK-NEXT: .type_name: int
783 ; CHECK-NEXT: .value_kind: by_value
784 ; CHECK-NEXT: - .name: b
785 ; CHECK-NEXT: .offset: 4
786 ; CHECK-NEXT: .size: 4
787 ; CHECK-NEXT: .type_name: short2
788 ; CHECK-NEXT: .value_kind: by_value
789 ; CHECK-NEXT: - .name: c
790 ; CHECK-NEXT: .offset: 8
791 ; CHECK-NEXT: .size: 4
792 ; CHECK-NEXT: .type_name: char3
793 ; CHECK-NEXT: .value_kind: by_value
794 ; CHECK-NEXT: - .offset: 16
795 ; CHECK-NEXT: .size: 8
796 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
797 ; CHECK-NEXT: - .offset: 24
798 ; CHECK-NEXT: .size: 8
799 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
800 ; CHECK-NEXT: - .offset: 32
801 ; CHECK-NEXT: .size: 8
802 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
803 ; CHECK-NEXT: - .address_space: global
804 ; CHECK-NEXT: .offset: 40
805 ; CHECK-NEXT: .size: 8
806 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
807 ; CHECK-NEXT: - .address_space: global
808 ; CHECK-NEXT: .offset: 48
809 ; CHECK-NEXT: .size: 8
810 ; CHECK-NEXT: .value_kind: hidden_none
811 ; CHECK-NEXT: - .address_space: global
812 ; CHECK-NEXT: .offset: 56
813 ; CHECK-NEXT: .size: 8
814 ; CHECK-NEXT: .value_kind: hidden_none
815 ; CHECK-NEXT: - .address_space: global
816 ; CHECK-NEXT: .offset: 64
817 ; CHECK-NEXT: .size: 8
818 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
819 ; CHECK: .language: OpenCL C
820 ; CHECK-NEXT: .language_version:
821 ; CHECK-NEXT: - 2
822 ; CHECK-NEXT: - 0
823 ; CHECK: .name: test_multi_arg
824 ; CHECK: .symbol: test_multi_arg.kd
831 ; CHECK: - .args:
832 ; CHECK-NEXT: - .address_space: global
833 ; CHECK-NEXT: .name: g
834 ; CHECK-NEXT: .offset: 0
835 ; CHECK-NEXT: .size: 8
836 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
837 ; CHECK-NEXT: .value_kind: global_buffer
838 ; CHECK-NEXT: - .address_space: constant
839 ; CHECK-NEXT: .name: c
840 ; CHECK-NEXT: .offset: 8
841 ; CHECK-NEXT: .size: 8
842 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
843 ; CHECK-NEXT: .value_kind: global_buffer
844 ; CHECK-NEXT: - .address_space: local
845 ; CHECK-NEXT: .name: l
846 ; CHECK-NEXT: .offset: 16
847 ; CHECK-NEXT: .pointee_align: 4
848 ; CHECK-NEXT: .size: 4
849 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
850 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
851 ; CHECK-NEXT: - .offset: 24
852 ; CHECK-NEXT: .size: 8
853 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
854 ; CHECK-NEXT: - .offset: 32
855 ; CHECK-NEXT: .size: 8
856 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
857 ; CHECK-NEXT: - .offset: 40
858 ; CHECK-NEXT: .size: 8
859 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
860 ; CHECK-NEXT: - .address_space: global
861 ; CHECK-NEXT: .offset: 48
862 ; CHECK-NEXT: .size: 8
863 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
864 ; CHECK-NEXT: - .address_space: global
865 ; CHECK-NEXT: .offset: 56
866 ; CHECK-NEXT: .size: 8
867 ; CHECK-NEXT: .value_kind: hidden_none
868 ; CHECK-NEXT: - .address_space: global
869 ; CHECK-NEXT: .offset: 64
870 ; CHECK-NEXT: .size: 8
871 ; CHECK-NEXT: .value_kind: hidden_none
872 ; CHECK-NEXT: - .address_space: global
873 ; CHECK-NEXT: .offset: 72
874 ; CHECK-NEXT: .size: 8
875 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
876 ; CHECK: .language: OpenCL C
877 ; CHECK-NEXT: .language_version:
878 ; CHECK-NEXT: - 2
879 ; CHECK-NEXT: - 0
880 ; CHECK: .name: test_addr_space
881 ; CHECK: .symbol: test_addr_space.kd
890 ; CHECK: - .args:
891 ; CHECK-NEXT: - .address_space: global
892 ; CHECK-NEXT: .is_volatile: true
893 ; CHECK-NEXT: .name: a
894 ; CHECK-NEXT: .offset: 0
895 ; CHECK-NEXT: .size: 8
896 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
897 ; CHECK-NEXT: .value_kind: global_buffer
898 ; CHECK-NEXT: - .address_space: global
899 ; CHECK-NEXT: .is_const: true
900 ; CHECK-NEXT: .is_restrict: true
901 ; CHECK-NEXT: .name: b
902 ; CHECK-NEXT: .offset: 8
903 ; CHECK-NEXT: .size: 8
904 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
905 ; CHECK-NEXT: .value_kind: global_buffer
906 ; CHECK-NEXT: - .address_space: global
907 ; CHECK-NEXT: .is_pipe: true
908 ; CHECK-NEXT: .name: c
909 ; CHECK-NEXT: .offset: 16
910 ; CHECK-NEXT: .size: 8
911 ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
912 ; CHECK-NEXT: .value_kind: pipe
913 ; CHECK-NEXT: - .offset: 24
914 ; CHECK-NEXT: .size: 8
915 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
916 ; CHECK-NEXT: - .offset: 32
917 ; CHECK-NEXT: .size: 8
918 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
919 ; CHECK-NEXT: - .offset: 40
920 ; CHECK-NEXT: .size: 8
921 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
922 ; CHECK-NEXT: - .address_space: global
923 ; CHECK-NEXT: .offset: 48
924 ; CHECK-NEXT: .size: 8
925 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
926 ; CHECK-NEXT: - .address_space: global
927 ; CHECK-NEXT: .offset: 56
928 ; CHECK-NEXT: .size: 8
929 ; CHECK-NEXT: .value_kind: hidden_none
930 ; CHECK-NEXT: - .address_space: global
931 ; CHECK-NEXT: .offset: 64
932 ; CHECK-NEXT: .size: 8
933 ; CHECK-NEXT: .value_kind: hidden_none
934 ; CHECK-NEXT: - .address_space: global
935 ; CHECK-NEXT: .offset: 72
936 ; CHECK-NEXT: .size: 8
937 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
938 ; CHECK: .language: OpenCL C
939 ; CHECK-NEXT: .language_version:
940 ; CHECK-NEXT: - 2
941 ; CHECK-NEXT: - 0
942 ; CHECK: .name: test_type_qual
943 ; CHECK: .symbol: test_type_qual.kd
952 ; CHECK: - .args:
953 ; CHECK-NEXT: - .access: read_only
954 ; CHECK-NEXT: .address_space: global
955 ; CHECK-NEXT: .name: ro
956 ; CHECK-NEXT: .offset: 0
957 ; CHECK-NEXT: .size: 8
958 ; CHECK-NEXT: .type_name: image1d_t
959 ; CHECK-NEXT: .value_kind: image
960 ; CHECK-NEXT: - .access: write_only
961 ; CHECK-NEXT: .address_space: global
962 ; CHECK-NEXT: .name: wo
963 ; CHECK-NEXT: .offset: 8
964 ; CHECK-NEXT: .size: 8
965 ; CHECK-NEXT: .type_name: image2d_t
966 ; CHECK-NEXT: .value_kind: image
967 ; CHECK-NEXT: - .access: read_write
968 ; CHECK-NEXT: .address_space: global
969 ; CHECK-NEXT: .name: rw
970 ; CHECK-NEXT: .offset: 16
971 ; CHECK-NEXT: .size: 8
972 ; CHECK-NEXT: .type_name: image3d_t
973 ; CHECK-NEXT: .value_kind: image
974 ; CHECK-NEXT: - .offset: 24
975 ; CHECK-NEXT: .size: 8
976 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
977 ; CHECK-NEXT: - .offset: 32
978 ; CHECK-NEXT: .size: 8
979 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
980 ; CHECK-NEXT: - .offset: 40
981 ; CHECK-NEXT: .size: 8
982 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
983 ; CHECK-NEXT: - .address_space: global
984 ; CHECK-NEXT: .offset: 48
985 ; CHECK-NEXT: .size: 8
986 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
987 ; CHECK-NEXT: - .address_space: global
988 ; CHECK-NEXT: .offset: 56
989 ; CHECK-NEXT: .size: 8
990 ; CHECK-NEXT: .value_kind: hidden_none
991 ; CHECK-NEXT: - .address_space: global
992 ; CHECK-NEXT: .offset: 64
993 ; CHECK-NEXT: .size: 8
994 ; CHECK-NEXT: .value_kind: hidden_none
995 ; CHECK-NEXT: - .address_space: global
996 ; CHECK-NEXT: .offset: 72
997 ; CHECK-NEXT: .size: 8
998 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
999 ; CHECK: .language: OpenCL C
1000 ; CHECK-NEXT: .language_version:
1001 ; CHECK-NEXT: - 2
1002 ; CHECK-NEXT: - 0
1003 ; CHECK: .name: test_access_qual
1004 ; CHECK: .symbol: test_access_qual.kd
1013 ; CHECK: - .args:
1014 ; CHECK-NEXT: - .name: a
1015 ; CHECK-NEXT: .offset: 0
1016 ; CHECK-NEXT: .size: 4
1017 ; CHECK-NEXT: .type_name: int
1018 ; CHECK-NEXT: .value_kind: by_value
1019 ; CHECK-NEXT: - .offset: 8
1020 ; CHECK-NEXT: .size: 8
1021 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1022 ; CHECK-NEXT: - .offset: 16
1023 ; CHECK-NEXT: .size: 8
1024 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1025 ; CHECK-NEXT: - .offset: 24
1026 ; CHECK-NEXT: .size: 8
1027 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1028 ; CHECK-NEXT: - .address_space: global
1029 ; CHECK-NEXT: .offset: 32
1030 ; CHECK-NEXT: .size: 8
1031 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1032 ; CHECK-NEXT: - .address_space: global
1033 ; CHECK-NEXT: .offset: 40
1034 ; CHECK-NEXT: .size: 8
1035 ; CHECK-NEXT: .value_kind: hidden_none
1036 ; CHECK-NEXT: - .address_space: global
1037 ; CHECK-NEXT: .offset: 48
1038 ; CHECK-NEXT: .size: 8
1039 ; CHECK-NEXT: .value_kind: hidden_none
1040 ; CHECK-NEXT: - .address_space: global
1041 ; CHECK-NEXT: .offset: 56
1042 ; CHECK-NEXT: .size: 8
1043 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1044 ; CHECK: .language: OpenCL C
1045 ; CHECK-NEXT: .language_version:
1046 ; CHECK-NEXT: - 2
1047 ; CHECK-NEXT: - 0
1048 ; CHECK: .name: test_vec_type_hint_half
1049 ; CHECK: .symbol: test_vec_type_hint_half.kd
1050 ; CHECK: .vec_type_hint: half
1057 ; CHECK: - .args:
1058 ; CHECK-NEXT: - .name: a
1059 ; CHECK-NEXT: .offset: 0
1060 ; CHECK-NEXT: .size: 4
1061 ; CHECK-NEXT: .type_name: int
1062 ; CHECK-NEXT: .value_kind: by_value
1063 ; CHECK-NEXT: - .offset: 8
1064 ; CHECK-NEXT: .size: 8
1065 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1066 ; CHECK-NEXT: - .offset: 16
1067 ; CHECK-NEXT: .size: 8
1068 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1069 ; CHECK-NEXT: - .offset: 24
1070 ; CHECK-NEXT: .size: 8
1071 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1072 ; CHECK-NEXT: - .address_space: global
1073 ; CHECK-NEXT: .offset: 32
1074 ; CHECK-NEXT: .size: 8
1075 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1076 ; CHECK-NEXT: - .address_space: global
1077 ; CHECK-NEXT: .offset: 40
1078 ; CHECK-NEXT: .size: 8
1079 ; CHECK-NEXT: .value_kind: hidden_none
1080 ; CHECK-NEXT: - .address_space: global
1081 ; CHECK-NEXT: .offset: 48
1082 ; CHECK-NEXT: .size: 8
1083 ; CHECK-NEXT: .value_kind: hidden_none
1084 ; CHECK-NEXT: - .address_space: global
1085 ; CHECK-NEXT: .offset: 56
1086 ; CHECK-NEXT: .size: 8
1087 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1088 ; CHECK: .language: OpenCL C
1089 ; CHECK-NEXT: .language_version:
1090 ; CHECK-NEXT: - 2
1091 ; CHECK-NEXT: - 0
1092 ; CHECK: .name: test_vec_type_hint_float
1093 ; CHECK: .symbol: test_vec_type_hint_float.kd
1094 ; CHECK: .vec_type_hint: float
1101 ; CHECK: - .args:
1102 ; CHECK-NEXT: - .name: a
1103 ; CHECK-NEXT: .offset: 0
1104 ; CHECK-NEXT: .size: 4
1105 ; CHECK-NEXT: .type_name: int
1106 ; CHECK-NEXT: .value_kind: by_value
1107 ; CHECK-NEXT: - .offset: 8
1108 ; CHECK-NEXT: .size: 8
1109 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1110 ; CHECK-NEXT: - .offset: 16
1111 ; CHECK-NEXT: .size: 8
1112 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1113 ; CHECK-NEXT: - .offset: 24
1114 ; CHECK-NEXT: .size: 8
1115 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1116 ; CHECK-NEXT: - .address_space: global
1117 ; CHECK-NEXT: .offset: 32
1118 ; CHECK-NEXT: .size: 8
1119 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1120 ; CHECK-NEXT: - .address_space: global
1121 ; CHECK-NEXT: .offset: 40
1122 ; CHECK-NEXT: .size: 8
1123 ; CHECK-NEXT: .value_kind: hidden_none
1124 ; CHECK-NEXT: - .address_space: global
1125 ; CHECK-NEXT: .offset: 48
1126 ; CHECK-NEXT: .size: 8
1127 ; CHECK-NEXT: .value_kind: hidden_none
1128 ; CHECK-NEXT: - .address_space: global
1129 ; CHECK-NEXT: .offset: 56
1130 ; CHECK-NEXT: .size: 8
1131 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1132 ; CHECK: .language: OpenCL C
1133 ; CHECK-NEXT: .language_version:
1134 ; CHECK-NEXT: - 2
1135 ; CHECK-NEXT: - 0
1136 ; CHECK: .name: test_vec_type_hint_double
1137 ; CHECK: .symbol: test_vec_type_hint_double.kd
1138 ; CHECK: .vec_type_hint: double
1145 ; CHECK: - .args:
1146 ; CHECK-NEXT: - .name: a
1147 ; CHECK-NEXT: .offset: 0
1148 ; CHECK-NEXT: .size: 4
1149 ; CHECK-NEXT: .type_name: int
1150 ; CHECK-NEXT: .value_kind: by_value
1151 ; CHECK-NEXT: - .offset: 8
1152 ; CHECK-NEXT: .size: 8
1153 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1154 ; CHECK-NEXT: - .offset: 16
1155 ; CHECK-NEXT: .size: 8
1156 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1157 ; CHECK-NEXT: - .offset: 24
1158 ; CHECK-NEXT: .size: 8
1159 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1160 ; CHECK-NEXT: - .address_space: global
1161 ; CHECK-NEXT: .offset: 32
1162 ; CHECK-NEXT: .size: 8
1163 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1164 ; CHECK-NEXT: - .address_space: global
1165 ; CHECK-NEXT: .offset: 40
1166 ; CHECK-NEXT: .size: 8
1167 ; CHECK-NEXT: .value_kind: hidden_none
1168 ; CHECK-NEXT: - .address_space: global
1169 ; CHECK-NEXT: .offset: 48
1170 ; CHECK-NEXT: .size: 8
1171 ; CHECK-NEXT: .value_kind: hidden_none
1172 ; CHECK-NEXT: - .address_space: global
1173 ; CHECK-NEXT: .offset: 56
1174 ; CHECK-NEXT: .size: 8
1175 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1176 ; CHECK: .language: OpenCL C
1177 ; CHECK-NEXT: .language_version:
1178 ; CHECK-NEXT: - 2
1179 ; CHECK-NEXT: - 0
1180 ; CHECK: .name: test_vec_type_hint_char
1181 ; CHECK: .symbol: test_vec_type_hint_char.kd
1182 ; CHECK: .vec_type_hint: char
1189 ; CHECK: - .args:
1190 ; CHECK-NEXT: - .name: a
1191 ; CHECK-NEXT: .offset: 0
1192 ; CHECK-NEXT: .size: 4
1193 ; CHECK-NEXT: .type_name: int
1194 ; CHECK-NEXT: .value_kind: by_value
1195 ; CHECK-NEXT: - .offset: 8
1196 ; CHECK-NEXT: .size: 8
1197 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1198 ; CHECK-NEXT: - .offset: 16
1199 ; CHECK-NEXT: .size: 8
1200 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1201 ; CHECK-NEXT: - .offset: 24
1202 ; CHECK-NEXT: .size: 8
1203 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1204 ; CHECK-NEXT: - .address_space: global
1205 ; CHECK-NEXT: .offset: 32
1206 ; CHECK-NEXT: .size: 8
1207 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1208 ; CHECK-NEXT: - .address_space: global
1209 ; CHECK-NEXT: .offset: 40
1210 ; CHECK-NEXT: .size: 8
1211 ; CHECK-NEXT: .value_kind: hidden_none
1212 ; CHECK-NEXT: - .address_space: global
1213 ; CHECK-NEXT: .offset: 48
1214 ; CHECK-NEXT: .size: 8
1215 ; CHECK-NEXT: .value_kind: hidden_none
1216 ; CHECK-NEXT: - .address_space: global
1217 ; CHECK-NEXT: .offset: 56
1218 ; CHECK-NEXT: .size: 8
1219 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1220 ; CHECK: .language: OpenCL C
1221 ; CHECK-NEXT: .language_version:
1222 ; CHECK-NEXT: - 2
1223 ; CHECK-NEXT: - 0
1224 ; CHECK: .name: test_vec_type_hint_short
1225 ; CHECK: .symbol: test_vec_type_hint_short.kd
1226 ; CHECK: .vec_type_hint: short
1233 ; CHECK: - .args:
1234 ; CHECK-NEXT: - .name: a
1235 ; CHECK-NEXT: .offset: 0
1236 ; CHECK-NEXT: .size: 4
1237 ; CHECK-NEXT: .type_name: int
1238 ; CHECK-NEXT: .value_kind: by_value
1239 ; CHECK-NEXT: - .offset: 8
1240 ; CHECK-NEXT: .size: 8
1241 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1242 ; CHECK-NEXT: - .offset: 16
1243 ; CHECK-NEXT: .size: 8
1244 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1245 ; CHECK-NEXT: - .offset: 24
1246 ; CHECK-NEXT: .size: 8
1247 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1248 ; CHECK-NEXT: - .address_space: global
1249 ; CHECK-NEXT: .offset: 32
1250 ; CHECK-NEXT: .size: 8
1251 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1252 ; CHECK-NEXT: - .address_space: global
1253 ; CHECK-NEXT: .offset: 40
1254 ; CHECK-NEXT: .size: 8
1255 ; CHECK-NEXT: .value_kind: hidden_none
1256 ; CHECK-NEXT: - .address_space: global
1257 ; CHECK-NEXT: .offset: 48
1258 ; CHECK-NEXT: .size: 8
1259 ; CHECK-NEXT: .value_kind: hidden_none
1260 ; CHECK-NEXT: - .address_space: global
1261 ; CHECK-NEXT: .offset: 56
1262 ; CHECK-NEXT: .size: 8
1263 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1264 ; CHECK: .language: OpenCL C
1265 ; CHECK-NEXT: .language_version:
1266 ; CHECK-NEXT: - 2
1267 ; CHECK-NEXT: - 0
1268 ; CHECK: .name: test_vec_type_hint_long
1269 ; CHECK: .symbol: test_vec_type_hint_long.kd
1270 ; CHECK: .vec_type_hint: long
1277 ; CHECK: - .args:
1278 ; CHECK-NEXT: - .name: a
1279 ; CHECK-NEXT: .offset: 0
1280 ; CHECK-NEXT: .size: 4
1281 ; CHECK-NEXT: .type_name: int
1282 ; CHECK-NEXT: .value_kind: by_value
1283 ; CHECK-NEXT: - .offset: 8
1284 ; CHECK-NEXT: .size: 8
1285 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1286 ; CHECK-NEXT: - .offset: 16
1287 ; CHECK-NEXT: .size: 8
1288 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1289 ; CHECK-NEXT: - .offset: 24
1290 ; CHECK-NEXT: .size: 8
1291 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1292 ; CHECK-NEXT: - .address_space: global
1293 ; CHECK-NEXT: .offset: 32
1294 ; CHECK-NEXT: .size: 8
1295 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1296 ; CHECK-NEXT: - .address_space: global
1297 ; CHECK-NEXT: .offset: 40
1298 ; CHECK-NEXT: .size: 8
1299 ; CHECK-NEXT: .value_kind: hidden_none
1300 ; CHECK-NEXT: - .address_space: global
1301 ; CHECK-NEXT: .offset: 48
1302 ; CHECK-NEXT: .size: 8
1303 ; CHECK-NEXT: .value_kind: hidden_none
1304 ; CHECK-NEXT: - .address_space: global
1305 ; CHECK-NEXT: .offset: 56
1306 ; CHECK-NEXT: .size: 8
1307 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1308 ; CHECK: .language: OpenCL C
1309 ; CHECK-NEXT: .language_version:
1310 ; CHECK-NEXT: - 2
1311 ; CHECK-NEXT: - 0
1312 ; CHECK: .name: test_vec_type_hint_unknown
1313 ; CHECK: .symbol: test_vec_type_hint_unknown.kd
1314 ; CHECK: .vec_type_hint: unknown
1321 ; CHECK: - .args:
1322 ; CHECK-NEXT: - .name: a
1323 ; CHECK-NEXT: .offset: 0
1324 ; CHECK-NEXT: .size: 4
1325 ; CHECK-NEXT: .type_name: int
1326 ; CHECK-NEXT: .value_kind: by_value
1327 ; CHECK-NEXT: - .offset: 8
1328 ; CHECK-NEXT: .size: 8
1329 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1330 ; CHECK-NEXT: - .offset: 16
1331 ; CHECK-NEXT: .size: 8
1332 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1333 ; CHECK-NEXT: - .offset: 24
1334 ; CHECK-NEXT: .size: 8
1335 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1336 ; CHECK-NEXT: - .address_space: global
1337 ; CHECK-NEXT: .offset: 32
1338 ; CHECK-NEXT: .size: 8
1339 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1340 ; CHECK-NEXT: - .address_space: global
1341 ; CHECK-NEXT: .offset: 40
1342 ; CHECK-NEXT: .size: 8
1343 ; CHECK-NEXT: .value_kind: hidden_none
1344 ; CHECK-NEXT: - .address_space: global
1345 ; CHECK-NEXT: .offset: 48
1346 ; CHECK-NEXT: .size: 8
1347 ; CHECK-NEXT: .value_kind: hidden_none
1348 ; CHECK-NEXT: - .address_space: global
1349 ; CHECK-NEXT: .offset: 56
1350 ; CHECK-NEXT: .size: 8
1351 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1352 ; CHECK: .language: OpenCL C
1353 ; CHECK-NEXT: .language_version:
1354 ; CHECK-NEXT: - 2
1355 ; CHECK-NEXT: - 0
1356 ; CHECK: .name: test_reqd_wgs_vec_type_hint
1357 ; CHECK: .reqd_workgroup_size:
1358 ; CHECK-NEXT: - 1
1359 ; CHECK-NEXT: - 2
1360 ; CHECK-NEXT: - 4
1361 ; CHECK: .symbol: test_reqd_wgs_vec_type_hint.kd
1362 ; CHECK: .vec_type_hint: int
1370 ; CHECK: - .args:
1371 ; CHECK-NEXT: - .name: a
1372 ; CHECK-NEXT: .offset: 0
1373 ; CHECK-NEXT: .size: 4
1374 ; CHECK-NEXT: .type_name: int
1375 ; CHECK-NEXT: .value_kind: by_value
1376 ; CHECK-NEXT: - .offset: 8
1377 ; CHECK-NEXT: .size: 8
1378 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1379 ; CHECK-NEXT: - .offset: 16
1380 ; CHECK-NEXT: .size: 8
1381 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1382 ; CHECK-NEXT: - .offset: 24
1383 ; CHECK-NEXT: .size: 8
1384 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1385 ; CHECK-NEXT: - .address_space: global
1386 ; CHECK-NEXT: .offset: 32
1387 ; CHECK-NEXT: .size: 8
1388 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1389 ; CHECK-NEXT: - .address_space: global
1390 ; CHECK-NEXT: .offset: 40
1391 ; CHECK-NEXT: .size: 8
1392 ; CHECK-NEXT: .value_kind: hidden_none
1393 ; CHECK-NEXT: - .address_space: global
1394 ; CHECK-NEXT: .offset: 48
1395 ; CHECK-NEXT: .size: 8
1396 ; CHECK-NEXT: .value_kind: hidden_none
1397 ; CHECK-NEXT: - .address_space: global
1398 ; CHECK-NEXT: .offset: 56
1399 ; CHECK-NEXT: .size: 8
1400 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1401 ; CHECK: .language: OpenCL C
1402 ; CHECK-NEXT: .language_version:
1403 ; CHECK-NEXT: - 2
1404 ; CHECK-NEXT: - 0
1405 ; CHECK: .name: test_wgs_hint_vec_type_hint
1406 ; CHECK: .symbol: test_wgs_hint_vec_type_hint.kd
1407 ; CHECK: .vec_type_hint: uint4
1408 ; CHECK: .workgroup_size_hint:
1409 ; CHECK-NEXT: - 8
1410 ; CHECK-NEXT: - 16
1411 ; CHECK-NEXT: - 32
1419 ; CHECK: - .args:
1420 ; CHECK-NEXT: - .address_space: global
1421 ; CHECK-NEXT: .name: a
1422 ; CHECK-NEXT: .offset: 0
1423 ; CHECK-NEXT: .size: 8
1424 ; CHECK-NEXT: .type_name: 'int addrspace(5)* addrspace(5)*'
1425 ; CHECK-NEXT: .value_kind: global_buffer
1426 ; CHECK-NEXT: - .offset: 8
1427 ; CHECK-NEXT: .size: 8
1428 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1429 ; CHECK-NEXT: - .offset: 16
1430 ; CHECK-NEXT: .size: 8
1431 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1432 ; CHECK-NEXT: - .offset: 24
1433 ; CHECK-NEXT: .size: 8
1434 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1435 ; CHECK-NEXT: - .address_space: global
1436 ; CHECK-NEXT: .offset: 32
1437 ; CHECK-NEXT: .size: 8
1438 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1439 ; CHECK-NEXT: - .address_space: global
1440 ; CHECK-NEXT: .offset: 40
1441 ; CHECK-NEXT: .size: 8
1442 ; CHECK-NEXT: .value_kind: hidden_none
1443 ; CHECK-NEXT: - .address_space: global
1444 ; CHECK-NEXT: .offset: 48
1445 ; CHECK-NEXT: .size: 8
1446 ; CHECK-NEXT: .value_kind: hidden_none
1447 ; CHECK-NEXT: - .address_space: global
1448 ; CHECK-NEXT: .offset: 56
1449 ; CHECK-NEXT: .size: 8
1450 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1451 ; CHECK: .language: OpenCL C
1452 ; CHECK-NEXT: .language_version:
1453 ; CHECK-NEXT: - 2
1454 ; CHECK-NEXT: - 0
1455 ; CHECK: .name: test_arg_ptr_to_ptr
1456 ; CHECK: .symbol: test_arg_ptr_to_ptr.kd
1463 ; CHECK: - .args:
1464 ; CHECK-NEXT: .name: a
1465 ; CHECK-NEXT: .offset: 0
1466 ; CHECK-NEXT: .size: 8
1467 ; CHECK-NEXT: .type_name: struct B
1468 ; CHECK-NEXT: .value_kind: by_value
1469 ; CHECK-NEXT: - .offset: 8
1470 ; CHECK-NEXT: .size: 8
1471 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1472 ; CHECK-NEXT: - .offset: 16
1473 ; CHECK-NEXT: .size: 8
1474 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1475 ; CHECK-NEXT: - .offset: 24
1476 ; CHECK-NEXT: .size: 8
1477 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1478 ; CHECK-NEXT: - .address_space: global
1479 ; CHECK-NEXT: .offset: 32
1480 ; CHECK-NEXT: .size: 8
1481 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1482 ; CHECK-NEXT: - .address_space: global
1483 ; CHECK-NEXT: .offset: 40
1484 ; CHECK-NEXT: .size: 8
1485 ; CHECK-NEXT: .value_kind: hidden_none
1486 ; CHECK-NEXT: - .address_space: global
1487 ; CHECK-NEXT: .offset: 48
1488 ; CHECK-NEXT: .size: 8
1489 ; CHECK-NEXT: .value_kind: hidden_none
1490 ; CHECK-NEXT: - .address_space: global
1491 ; CHECK-NEXT: .offset: 56
1492 ; CHECK-NEXT: .size: 8
1493 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1494 ; CHECK: .language: OpenCL C
1495 ; CHECK-NEXT: .language_version:
1496 ; CHECK-NEXT: - 2
1497 ; CHECK-NEXT: - 0
1498 ; CHECK: .name: test_arg_struct_contains_ptr
1499 ; CHECK: .symbol: test_arg_struct_contains_ptr.kd
1506 ; CHECK: - .args:
1507 ; CHECK-NEXT: - .name: a
1508 ; CHECK-NEXT: .offset: 0
1509 ; CHECK-NEXT: .size: 16
1510 ; CHECK-NEXT: .type_name: 'global int addrspace(5)* __attribute__((ext_vector_type(2))…
1511 ; CHECK-NEXT: .value_kind: by_value
1512 ; CHECK-NEXT: - .offset: 16
1513 ; CHECK-NEXT: .size: 8
1514 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1515 ; CHECK-NEXT: - .offset: 24
1516 ; CHECK-NEXT: .size: 8
1517 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1518 ; CHECK-NEXT: - .offset: 32
1519 ; CHECK-NEXT: .size: 8
1520 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1521 ; CHECK-NEXT: - .address_space: global
1522 ; CHECK-NEXT: .offset: 40
1523 ; CHECK-NEXT: .size: 8
1524 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1525 ; CHECK-NEXT: - .address_space: global
1526 ; CHECK-NEXT: .offset: 48
1527 ; CHECK-NEXT: .size: 8
1528 ; CHECK-NEXT: .value_kind: hidden_none
1529 ; CHECK-NEXT: - .address_space: global
1530 ; CHECK-NEXT: .offset: 56
1531 ; CHECK-NEXT: .size: 8
1532 ; CHECK-NEXT: .value_kind: hidden_none
1533 ; CHECK-NEXT: - .address_space: global
1534 ; CHECK-NEXT: .offset: 64
1535 ; CHECK-NEXT: .size: 8
1536 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1537 ; CHECK: .language: OpenCL C
1538 ; CHECK-NEXT: .language_version:
1539 ; CHECK-NEXT: - 2
1540 ; CHECK-NEXT: - 0
1541 ; CHECK: .name: test_arg_vector_of_ptr
1542 ; CHECK: .symbol: test_arg_vector_of_ptr.kd
1549 ; CHECK: - .args:
1550 ; CHECK-NEXT: - .address_space: global
1551 ; CHECK-NEXT: .name: a
1552 ; CHECK-NEXT: .offset: 0
1553 ; CHECK-NEXT: .size: 8
1554 ; CHECK-NEXT: .type_name: clk_event_t
1555 ; CHECK-NEXT: .value_kind: global_buffer
1556 ; CHECK-NEXT: - .offset: 8
1557 ; CHECK-NEXT: .size: 8
1558 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1559 ; CHECK-NEXT: - .offset: 16
1560 ; CHECK-NEXT: .size: 8
1561 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1562 ; CHECK-NEXT: - .offset: 24
1563 ; CHECK-NEXT: .size: 8
1564 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1565 ; CHECK-NEXT: - .address_space: global
1566 ; CHECK-NEXT: .offset: 32
1567 ; CHECK-NEXT: .size: 8
1568 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1569 ; CHECK-NEXT: - .address_space: global
1570 ; CHECK-NEXT: .offset: 40
1571 ; CHECK-NEXT: .size: 8
1572 ; CHECK-NEXT: .value_kind: hidden_none
1573 ; CHECK-NEXT: - .address_space: global
1574 ; CHECK-NEXT: .offset: 48
1575 ; CHECK-NEXT: .size: 8
1576 ; CHECK-NEXT: .value_kind: hidden_none
1577 ; CHECK-NEXT: - .address_space: global
1578 ; CHECK-NEXT: .offset: 56
1579 ; CHECK-NEXT: .size: 8
1580 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1581 ; CHECK: .language: OpenCL C
1582 ; CHECK-NEXT: .language_version:
1583 ; CHECK-NEXT: - 2
1584 ; CHECK-NEXT: - 0
1585 ; CHECK: .name: test_arg_unknown_builtin_type
1586 ; CHECK: .symbol: test_arg_unknown_builtin_type.kd
1594 ; CHECK: - .args:
1595 ; CHECK-NEXT: - .address_space: global
1596 ; CHECK-NEXT: .name: a
1597 ; CHECK-NEXT: .offset: 0
1598 ; CHECK-NEXT: .size: 8
1599 ; CHECK-NEXT: .type_name: 'long addrspace(5)*'
1600 ; CHECK-NEXT: .value_kind: global_buffer
1601 ; CHECK-NEXT: - .address_space: local
1602 ; CHECK-NEXT: .name: b
1603 ; CHECK-NEXT: .offset: 8
1604 ; CHECK-NEXT: .pointee_align: 1
1605 ; CHECK-NEXT: .size: 4
1606 ; CHECK-NEXT: .type_name: 'char addrspace(5)*'
1607 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1608 ; CHECK-NEXT: - .address_space: local
1609 ; CHECK-NEXT: .name: c
1610 ; CHECK-NEXT: .offset: 12
1611 ; CHECK-NEXT: .pointee_align: 2
1612 ; CHECK-NEXT: .size: 4
1613 ; CHECK-NEXT: .type_name: 'char2 addrspace(5)*'
1614 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1615 ; CHECK-NEXT: - .address_space: local
1616 ; CHECK-NEXT: .name: d
1617 ; CHECK-NEXT: .offset: 16
1618 ; CHECK-NEXT: .pointee_align: 4
1619 ; CHECK-NEXT: .size: 4
1620 ; CHECK-NEXT: .type_name: 'char3 addrspace(5)*'
1621 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1622 ; CHECK-NEXT: - .address_space: local
1623 ; CHECK-NEXT: .name: e
1624 ; CHECK-NEXT: .offset: 20
1625 ; CHECK-NEXT: .pointee_align: 4
1626 ; CHECK-NEXT: .size: 4
1627 ; CHECK-NEXT: .type_name: 'char4 addrspace(5)*'
1628 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1629 ; CHECK-NEXT: - .address_space: local
1630 ; CHECK-NEXT: .name: f
1631 ; CHECK-NEXT: .offset: 24
1632 ; CHECK-NEXT: .pointee_align: 8
1633 ; CHECK-NEXT: .size: 4
1634 ; CHECK-NEXT: .type_name: 'char8 addrspace(5)*'
1635 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1636 ; CHECK-NEXT: - .address_space: local
1637 ; CHECK-NEXT: .name: g
1638 ; CHECK-NEXT: .offset: 28
1639 ; CHECK-NEXT: .pointee_align: 16
1640 ; CHECK-NEXT: .size: 4
1641 ; CHECK-NEXT: .type_name: 'char16 addrspace(5)*'
1642 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1643 ; CHECK-NEXT: - .address_space: local
1644 ; CHECK-NEXT: .name: h
1645 ; CHECK-NEXT: .offset: 32
1646 ; CHECK-NEXT: .pointee_align: 1
1647 ; CHECK-NEXT: .size: 4
1648 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1649 ; CHECK-NEXT: - .offset: 40
1650 ; CHECK-NEXT: .size: 8
1651 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1652 ; CHECK-NEXT: - .offset: 48
1653 ; CHECK-NEXT: .size: 8
1654 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1655 ; CHECK-NEXT: - .offset: 56
1656 ; CHECK-NEXT: .size: 8
1657 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1658 ; CHECK-NEXT: - .address_space: global
1659 ; CHECK-NEXT: .offset: 64
1660 ; CHECK-NEXT: .size: 8
1661 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1662 ; CHECK-NEXT: - .address_space: global
1663 ; CHECK-NEXT: .offset: 72
1664 ; CHECK-NEXT: .size: 8
1665 ; CHECK-NEXT: .value_kind: hidden_none
1666 ; CHECK-NEXT: - .address_space: global
1667 ; CHECK-NEXT: .offset: 80
1668 ; CHECK-NEXT: .size: 8
1669 ; CHECK-NEXT: .value_kind: hidden_none
1670 ; CHECK-NEXT: - .address_space: global
1671 ; CHECK-NEXT: .offset: 88
1672 ; CHECK-NEXT: .size: 8
1673 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1674 ; CHECK: .language: OpenCL C
1675 ; CHECK-NEXT: .language_version:
1676 ; CHECK-NEXT: - 2
1677 ; CHECK-NEXT: - 0
1678 ; CHECK: .name: test_pointee_align
1679 ; CHECK: .symbol: test_pointee_align.kd
1693 ; CHECK: - .args:
1694 ; CHECK-NEXT: - .address_space: global
1695 ; CHECK-NEXT: .name: a
1696 ; CHECK-NEXT: .offset: 0
1697 ; CHECK-NEXT: .size: 8
1698 ; CHECK-NEXT: .type_name: 'long addrspace(5)*'
1699 ; CHECK-NEXT: .value_kind: global_buffer
1700 ; CHECK-NEXT: - .address_space: local
1701 ; CHECK-NEXT: .name: b
1702 ; CHECK-NEXT: .offset: 8
1703 ; CHECK-NEXT: .pointee_align: 8
1704 ; CHECK-NEXT: .size: 4
1705 ; CHECK-NEXT: .type_name: 'char addrspace(5)*'
1706 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1707 ; CHECK-NEXT: - .address_space: local
1708 ; CHECK-NEXT: .name: c
1709 ; CHECK-NEXT: .offset: 12
1710 ; CHECK-NEXT: .pointee_align: 32
1711 ; CHECK-NEXT: .size: 4
1712 ; CHECK-NEXT: .type_name: 'char2 addrspace(5)*'
1713 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1714 ; CHECK-NEXT: - .address_space: local
1715 ; CHECK-NEXT: .name: d
1716 ; CHECK-NEXT: .offset: 16
1717 ; CHECK-NEXT: .pointee_align: 64
1718 ; CHECK-NEXT: .size: 4
1719 ; CHECK-NEXT: .type_name: 'char3 addrspace(5)*'
1720 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1721 ; CHECK-NEXT: - .address_space: local
1722 ; CHECK-NEXT: .name: e
1723 ; CHECK-NEXT: .offset: 20
1724 ; CHECK-NEXT: .pointee_align: 256
1725 ; CHECK-NEXT: .size: 4
1726 ; CHECK-NEXT: .type_name: 'char4 addrspace(5)*'
1727 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1728 ; CHECK-NEXT: - .address_space: local
1729 ; CHECK-NEXT: .name: f
1730 ; CHECK-NEXT: .offset: 24
1731 ; CHECK-NEXT: .pointee_align: 128
1732 ; CHECK-NEXT: .size: 4
1733 ; CHECK-NEXT: .type_name: 'char8 addrspace(5)*'
1734 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1735 ; CHECK-NEXT: - .address_space: local
1736 ; CHECK-NEXT: .name: g
1737 ; CHECK-NEXT: .offset: 28
1738 ; CHECK-NEXT: .pointee_align: 1024
1739 ; CHECK-NEXT: .size: 4
1740 ; CHECK-NEXT: .type_name: 'char16 addrspace(5)*'
1741 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1742 ; CHECK-NEXT: - .address_space: local
1743 ; CHECK-NEXT: .name: h
1744 ; CHECK-NEXT: .offset: 32
1745 ; CHECK-NEXT: .pointee_align: 16
1746 ; CHECK-NEXT: .size: 4
1747 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1748 ; CHECK-NEXT: - .offset: 40
1749 ; CHECK-NEXT: .size: 8
1750 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1751 ; CHECK-NEXT: - .offset: 48
1752 ; CHECK-NEXT: .size: 8
1753 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1754 ; CHECK-NEXT: - .offset: 56
1755 ; CHECK-NEXT: .size: 8
1756 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1757 ; CHECK-NEXT: - .address_space: global
1758 ; CHECK-NEXT: .offset: 64
1759 ; CHECK-NEXT: .size: 8
1760 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1761 ; CHECK-NEXT: - .address_space: global
1762 ; CHECK-NEXT: .offset: 72
1763 ; CHECK-NEXT: .size: 8
1764 ; CHECK-NEXT: .value_kind: hidden_none
1765 ; CHECK-NEXT: - .address_space: global
1766 ; CHECK-NEXT: .offset: 80
1767 ; CHECK-NEXT: .size: 8
1768 ; CHECK-NEXT: .value_kind: hidden_none
1769 ; CHECK-NEXT: - .address_space: global
1770 ; CHECK-NEXT: .offset: 88
1771 ; CHECK-NEXT: .size: 8
1772 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1773 ; CHECK: .language: OpenCL C
1774 ; CHECK-NEXT: .language_version:
1775 ; CHECK-NEXT: - 2
1776 ; CHECK-NEXT: - 0
1777 ; CHECK: .name: test_pointee_align_attribute
1778 ; CHECK: .symbol: test_pointee_align_attribute.kd
1791 ; CHECK: - .args:
1792 ; CHECK-NEXT: - .name: arg
1793 ; CHECK-NEXT: .offset: 0
1794 ; CHECK-NEXT: .size: 25
1795 ; CHECK-NEXT: .type_name: __block_literal
1796 ; CHECK-NEXT: .value_kind: by_value
1797 ; CHECK-NEXT: - .offset: 32
1798 ; CHECK-NEXT: .size: 8
1799 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1800 ; CHECK-NEXT: - .offset: 40
1801 ; CHECK-NEXT: .size: 8
1802 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1803 ; CHECK-NEXT: - .offset: 48
1804 ; CHECK-NEXT: .size: 8
1805 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1806 ; CHECK-NEXT: - .address_space: global
1807 ; CHECK-NEXT: .offset: 56
1808 ; CHECK-NEXT: .size: 8
1809 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1810 ; CHECK-NEXT: - .address_space: global
1811 ; CHECK-NEXT: .offset: 64
1812 ; CHECK-NEXT: .size: 8
1813 ; CHECK-NEXT: .value_kind: hidden_none
1814 ; CHECK-NEXT: - .address_space: global
1815 ; CHECK-NEXT: .offset: 72
1816 ; CHECK-NEXT: .size: 8
1817 ; CHECK-NEXT: .value_kind: hidden_none
1818 ; CHECK-NEXT: - .address_space: global
1819 ; CHECK-NEXT: .offset: 80
1820 ; CHECK-NEXT: .size: 8
1821 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1822 ; CHECK: .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle
1823 ; CHECK: .language: OpenCL C
1824 ; CHECK-NEXT: .language_version:
1825 ; CHECK-NEXT: - 2
1826 ; CHECK-NEXT: - 0
1827 ; CHECK: .name: __test_block_invoke_kernel
1828 ; CHECK: .symbol: __test_block_invoke_kernel.kd
1836 ; CHECK: - .args:
1837 ; CHECK-NEXT: - .name: a
1838 ; CHECK-NEXT: .offset: 0
1839 ; CHECK-NEXT: .size: 1
1840 ; CHECK-NEXT: .type_name: char
1841 ; CHECK-NEXT: .value_kind: by_value
1842 ; CHECK-NEXT: - .offset: 8
1843 ; CHECK-NEXT: .size: 8
1844 ; CHECK-NEXT: .value_kind: hidden_global_offset_x
1845 ; CHECK-NEXT: - .offset: 16
1846 ; CHECK-NEXT: .size: 8
1847 ; CHECK-NEXT: .value_kind: hidden_global_offset_y
1848 ; CHECK-NEXT: - .offset: 24
1849 ; CHECK-NEXT: .size: 8
1850 ; CHECK-NEXT: .value_kind: hidden_global_offset_z
1851 ; CHECK-NEXT: - .address_space: global
1852 ; CHECK-NEXT: .offset: 32
1853 ; CHECK-NEXT: .size: 8
1854 ; CHECK-NEXT: .value_kind: hidden_printf_buffer
1855 ; CHECK-NEXT: - .address_space: global
1856 ; CHECK-NEXT: .offset: 40
1857 ; CHECK-NEXT: .size: 8
1858 ; CHECK-NEXT: .value_kind: hidden_default_queue
1859 ; CHECK-NEXT: - .address_space: global
1860 ; CHECK-NEXT: .offset: 48
1861 ; CHECK-NEXT: .size: 8
1862 ; CHECK-NEXT: .value_kind: hidden_completion_action
1863 ; CHECK-NEXT: - .address_space: global
1864 ; CHECK-NEXT: .offset: 56
1865 ; CHECK-NEXT: .size: 8
1866 ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
1867 ; CHECK: .language: OpenCL C
1868 ; CHECK-NEXT: .language_version:
1869 ; CHECK-NEXT: - 2
1870 ; CHECK-NEXT: - 0
1871 ; CHECK: .name: test_enqueue_kernel_caller
1872 ; CHECK: .symbol: test_enqueue_kernel_caller.kd
1879 ; CHECK: - .args:
1880 ; CHECK-NEXT: - .name: ptr
1881 ; CHECK-NEXT: .offset: 0
1882 ; CHECK-NEXT: .size: 8
1883 ; CHECK-NEXT: .value_kind: global_buffer
1884 ; CHECK: .name: unknown_addrspace_kernarg
1885 ; CHECK: .symbol: unknown_addrspace_kernarg.kd
1890 ; CHECK: amdhsa.printf:
1891 ; CHECK-NEXT: - '1:1:4:%d\n'
1892 ; CHECK-NEXT: - '2:1:8:%g\n'
1893 ; CHECK: amdhsa.version:
1894 ; CHECK-NEXT: - 1
1895 ; CHECK-NEXT: - 0