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