| blob | 85ee9b4858b4defa898193ecc5835518ebb1cab1 |
1 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX700 --check-prefix=NOTES %s
2 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx802 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX802 --check-prefix=NOTES %s
3 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX900 --check-prefix=NOTES %s
4 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
5 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx802 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
6 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
8 %struct.A = type { i8, float }
9 %opencl.image1d_t = type opaque
10 %opencl.image2d_t = type opaque
11 %opencl.image3d_t = type opaque
12 %opencl.queue_t = type opaque
13 %opencl.pipe_t = type opaque
14 %struct.B = type { 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: amdhsa.kernels:
21 ; CHECK: .symbol: test_char.kd
22 ; CHECK: .name: test_char
23 ; CHECK: .language: OpenCL C
24 ; CHECK: .language_version:
25 ; CHECK-NEXT: - 2
26 ; CHECK-NEXT: - 0
27 ; CHECK: .args:
28 ; CHECK-NEXT: - .type_name: char
29 ; CHECK-NEXT: .value_kind: by_value
30 ; CHECK-NEXT: .offset: 0
31 ; CHECK-NEXT: .size: 1
32 ; CHECK-NEXT: .value_type: i8
33 ; CHECK-NEXT: .name: a
34 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
35 ; CHECK-NEXT: .offset: 8
36 ; CHECK-NEXT: .size: 8
37 ; CHECK-NEXT: .value_type: i64
38 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
39 ; CHECK-NEXT: .offset: 16
40 ; CHECK-NEXT: .size: 8
41 ; CHECK-NEXT: .value_type: i64
42 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
43 ; CHECK-NEXT: .offset: 24
44 ; CHECK-NEXT: .size: 8
45 ; CHECK-NEXT: .value_type: i64
46 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
47 ; CHECK-NEXT: .offset: 32
48 ; CHECK-NEXT: .size: 8
49 ; CHECK-NEXT: .value_type: i8
50 ; CHECK-NEXT: .address_space: global
51 ; CHECK-NOT: .value_kind: hidden_default_queue
52 ; CHECK-NOT: .value_kind: hidden_completion_action
53 define amdgpu_kernel void @test_char(i8 %a)
54 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
55 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
56 ret void
57 }
59 ; CHECK: .symbol: test_ushort2.kd
60 ; CHECK: .name: test_ushort2
61 ; CHECK: .language: OpenCL C
62 ; CHECK: .language_version:
63 ; CHECK-NEXT: - 2
64 ; CHECK-NEXT: - 0
65 ; CHECK: .args:
66 ; CHECK-NEXT: - .type_name: ushort2
67 ; CHECK-NEXT: .value_kind: by_value
68 ; CHECK-NEXT: .offset: 0
69 ; CHECK-NEXT: .size: 4
70 ; CHECK-NEXT: .value_type: u16
71 ; CHECK-NEXT: .name: a
72 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
73 ; CHECK-NEXT: .offset: 8
74 ; CHECK-NEXT: .size: 8
75 ; CHECK-NEXT: .value_type: i64
76 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
77 ; CHECK-NEXT: .offset: 16
78 ; CHECK-NEXT: .size: 8
79 ; CHECK-NEXT: .value_type: i64
80 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
81 ; CHECK-NEXT: .offset: 24
82 ; CHECK-NEXT: .size: 8
83 ; CHECK-NEXT: .value_type: i64
84 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
85 ; CHECK-NEXT: .offset: 32
86 ; CHECK-NEXT: .size: 8
87 ; CHECK-NEXT: .value_type: i8
88 ; CHECK-NEXT: .address_space: global
89 define amdgpu_kernel void @test_ushort2(<2 x i16> %a)
90 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !10
91 !kernel_arg_base_type !10 !kernel_arg_type_qual !4 {
92 ret void
93 }
95 ; CHECK: .symbol: test_int3.kd
96 ; CHECK: .name: test_int3
97 ; CHECK: .language: OpenCL C
98 ; CHECK: .language_version:
99 ; CHECK-NEXT: - 2
100 ; CHECK-NEXT: - 0
101 ; CHECK: .args:
102 ; CHECK-NEXT: - .type_name: int3
103 ; CHECK-NEXT: .value_kind: by_value
104 ; CHECK-NEXT: .offset: 0
105 ; CHECK-NEXT: .size: 16
106 ; CHECK-NEXT: .value_type: i32
107 ; CHECK-NEXT: .name: a
108 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
109 ; CHECK-NEXT: .offset: 16
110 ; CHECK-NEXT: .size: 8
111 ; CHECK-NEXT: .value_type: i64
112 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
113 ; CHECK-NEXT: .offset: 24
114 ; CHECK-NEXT: .size: 8
115 ; CHECK-NEXT: .value_type: i64
116 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
117 ; CHECK-NEXT: .offset: 32
118 ; CHECK-NEXT: .size: 8
119 ; CHECK-NEXT: .value_type: i64
120 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
121 ; CHECK-NEXT: .offset: 40
122 ; CHECK-NEXT: .size: 8
123 ; CHECK-NEXT: .value_type: i8
124 ; CHECK-NEXT: .address_space: global
125 define amdgpu_kernel void @test_int3(<3 x i32> %a)
126 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !11
127 !kernel_arg_base_type !11 !kernel_arg_type_qual !4 {
128 ret void
129 }
131 ; CHECK: .symbol: test_ulong4.kd
132 ; CHECK: .name: test_ulong4
133 ; CHECK: .language: OpenCL C
134 ; CHECK: .language_version:
135 ; CHECK-NEXT: - 2
136 ; CHECK-NEXT: - 0
137 ; CHECK: .args:
138 ; CHECK-NEXT: - .type_name: ulong4
139 ; CHECK-NEXT: .value_kind: by_value
140 ; CHECK-NEXT: .offset: 0
141 ; CHECK-NEXT: .size: 32
142 ; CHECK-NEXT: .value_type: u64
143 ; CHECK-NEXT: .name: a
144 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
145 ; CHECK-NEXT: .offset: 32
146 ; CHECK-NEXT: .size: 8
147 ; CHECK-NEXT: .value_type: i64
148 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
149 ; CHECK-NEXT: .offset: 40
150 ; CHECK-NEXT: .size: 8
151 ; CHECK-NEXT: .value_type: i64
152 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
153 ; CHECK-NEXT: .offset: 48
154 ; CHECK-NEXT: .size: 8
155 ; CHECK-NEXT: .value_type: i64
156 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
157 ; CHECK-NEXT: .offset: 56
158 ; CHECK-NEXT: .size: 8
159 ; CHECK-NEXT: .value_type: i8
160 ; CHECK-NEXT: .address_space: global
161 define amdgpu_kernel void @test_ulong4(<4 x i64> %a)
162 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !12
163 !kernel_arg_base_type !12 !kernel_arg_type_qual !4 {
164 ret void
165 }
167 ; CHECK: .symbol: test_half8.kd
168 ; CHECK: .name: test_half8
169 ; CHECK: .language: OpenCL C
170 ; CHECK: .language_version:
171 ; CHECK-NEXT: - 2
172 ; CHECK-NEXT: - 0
173 ; CHECK: .args:
174 ; CHECK-NEXT: - .type_name: half8
175 ; CHECK-NEXT: .value_kind: by_value
176 ; CHECK-NEXT: .offset: 0
177 ; CHECK-NEXT: .size: 16
178 ; CHECK-NEXT: .value_type: f16
179 ; CHECK-NEXT: .name: a
180 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
181 ; CHECK-NEXT: .offset: 16
182 ; CHECK-NEXT: .size: 8
183 ; CHECK-NEXT: .value_type: i64
184 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
185 ; CHECK-NEXT: .offset: 24
186 ; CHECK-NEXT: .size: 8
187 ; CHECK-NEXT: .value_type: i64
188 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
189 ; CHECK-NEXT: .offset: 32
190 ; CHECK-NEXT: .size: 8
191 ; CHECK-NEXT: .value_type: i64
192 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
193 ; CHECK-NEXT: .offset: 40
194 ; CHECK-NEXT: .size: 8
195 ; CHECK-NEXT: .value_type: i8
196 ; CHECK-NEXT: .address_space: global
197 define amdgpu_kernel void @test_half8(<8 x half> %a)
198 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !13
199 !kernel_arg_base_type !13 !kernel_arg_type_qual !4 {
200 ret void
201 }
203 ; CHECK: .symbol: test_float16.kd
204 ; CHECK: .name: test_float16
205 ; CHECK: .language: OpenCL C
206 ; CHECK: .language_version:
207 ; CHECK-NEXT: - 2
208 ; CHECK-NEXT: - 0
209 ; CHECK: .args:
210 ; CHECK-NEXT: - .type_name: float16
211 ; CHECK-NEXT: .value_kind: by_value
212 ; CHECK-NEXT: .offset: 0
213 ; CHECK-NEXT: .size: 64
214 ; CHECK-NEXT: .value_type: f32
215 ; CHECK-NEXT: .name: a
216 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
217 ; CHECK-NEXT: .offset: 64
218 ; CHECK-NEXT: .size: 8
219 ; CHECK-NEXT: .value_type: i64
220 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
221 ; CHECK-NEXT: .offset: 72
222 ; CHECK-NEXT: .size: 8
223 ; CHECK-NEXT: .value_type: i64
224 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
225 ; CHECK-NEXT: .offset: 80
226 ; CHECK-NEXT: .size: 8
227 ; CHECK-NEXT: .value_type: i64
228 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
229 ; CHECK-NEXT: .offset: 88
230 ; CHECK-NEXT: .size: 8
231 ; CHECK-NEXT: .value_type: i8
232 ; CHECK-NEXT: .address_space: global
233 define amdgpu_kernel void @test_float16(<16 x float> %a)
234 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !14
235 !kernel_arg_base_type !14 !kernel_arg_type_qual !4 {
236 ret void
237 }
239 ; CHECK: .symbol: test_double16.kd
240 ; CHECK: .name: test_double16
241 ; CHECK: .language: OpenCL C
242 ; CHECK: .language_version:
243 ; CHECK-NEXT: - 2
244 ; CHECK-NEXT: - 0
245 ; CHECK: .args:
246 ; CHECK-NEXT: - .type_name: double16
247 ; CHECK-NEXT: .value_kind: by_value
248 ; CHECK-NEXT: .offset: 0
249 ; CHECK-NEXT: .size: 128
250 ; CHECK-NEXT: .value_type: f64
251 ; CHECK-NEXT: .name: a
252 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
253 ; CHECK-NEXT: .offset: 128
254 ; CHECK-NEXT: .size: 8
255 ; CHECK-NEXT: .value_type: i64
256 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
257 ; CHECK-NEXT: .offset: 136
258 ; CHECK-NEXT: .size: 8
259 ; CHECK-NEXT: .value_type: i64
260 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
261 ; CHECK-NEXT: .offset: 144
262 ; CHECK-NEXT: .size: 8
263 ; CHECK-NEXT: .value_type: i64
264 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
265 ; CHECK-NEXT: .offset: 152
266 ; CHECK-NEXT: .size: 8
267 ; CHECK-NEXT: .value_type: i8
268 ; CHECK-NEXT: .address_space: global
269 define amdgpu_kernel void @test_double16(<16 x double> %a)
270 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !15
271 !kernel_arg_base_type !15 !kernel_arg_type_qual !4 {
272 ret void
273 }
275 ; CHECK: .symbol: test_pointer.kd
276 ; CHECK: .name: test_pointer
277 ; CHECK: .language: OpenCL C
278 ; CHECK: .language_version:
279 ; CHECK-NEXT: - 2
280 ; CHECK-NEXT: - 0
281 ; CHECK: .args:
282 ; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
283 ; CHECK-NEXT: .value_kind: global_buffer
284 ; CHECK-NEXT: .name: a
285 ; CHECK-NEXT: .offset: 0
286 ; CHECK-NEXT: .size: 8
287 ; CHECK-NEXT: .value_type: i32
288 ; CHECK-NEXT: .address_space: global
289 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
290 ; CHECK-NEXT: .offset: 8
291 ; CHECK-NEXT: .size: 8
292 ; CHECK-NEXT: .value_type: i64
293 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
294 ; CHECK-NEXT: .offset: 16
295 ; CHECK-NEXT: .size: 8
296 ; CHECK-NEXT: .value_type: i64
297 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
298 ; CHECK-NEXT: .offset: 24
299 ; CHECK-NEXT: .size: 8
300 ; CHECK-NEXT: .value_type: i64
301 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
302 ; CHECK-NEXT: .offset: 32
303 ; CHECK-NEXT: .size: 8
304 ; CHECK-NEXT: .value_type: i8
305 ; CHECK-NEXT: .address_space: global
306 define amdgpu_kernel void @test_pointer(i32 addrspace(1)* %a)
307 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !16
308 !kernel_arg_base_type !16 !kernel_arg_type_qual !4 {
309 ret void
310 }
312 ; CHECK: .symbol: test_image.kd
313 ; CHECK: .name: test_image
314 ; CHECK: .language: OpenCL C
315 ; CHECK: .language_version:
316 ; CHECK-NEXT: - 2
317 ; CHECK-NEXT: - 0
318 ; CHECK: .args:
319 ; CHECK-NEXT: - .type_name: image2d_t
320 ; CHECK-NEXT: .value_kind: image
321 ; CHECK-NEXT: .name: a
322 ; CHECK-NEXT: .offset: 0
323 ; CHECK-NEXT: .size: 8
324 ; CHECK-NEXT: .value_type: struct
325 ; CHECK-NEXT: .address_space: global
326 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
327 ; CHECK-NEXT: .offset: 8
328 ; CHECK-NEXT: .size: 8
329 ; CHECK-NEXT: .value_type: i64
330 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
331 ; CHECK-NEXT: .offset: 16
332 ; CHECK-NEXT: .size: 8
333 ; CHECK-NEXT: .value_type: i64
334 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
335 ; CHECK-NEXT: .offset: 24
336 ; CHECK-NEXT: .size: 8
337 ; CHECK-NEXT: .value_type: i64
338 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
339 ; CHECK-NEXT: .offset: 32
340 ; CHECK-NEXT: .size: 8
341 ; CHECK-NEXT: .value_type: i8
342 ; CHECK-NEXT: .address_space: global
343 define amdgpu_kernel void @test_image(%opencl.image2d_t addrspace(1)* %a)
344 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !17
345 !kernel_arg_base_type !17 !kernel_arg_type_qual !4 {
346 ret void
347 }
349 ; CHECK: .symbol: test_sampler.kd
350 ; CHECK: .name: test_sampler
351 ; CHECK: .language: OpenCL C
352 ; CHECK: .language_version:
353 ; CHECK-NEXT: - 2
354 ; CHECK-NEXT: - 0
355 ; CHECK: .args:
356 ; CHECK-NEXT: - .type_name: sampler_t
357 ; CHECK-NEXT: .value_kind: sampler
358 ; CHECK-NEXT: .offset: 0
359 ; CHECK-NEXT: .size: 4
360 ; CHECK-NEXT: .value_type: i32
361 ; CHECK-NEXT: .name: a
362 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
363 ; CHECK-NEXT: .offset: 8
364 ; CHECK-NEXT: .size: 8
365 ; CHECK-NEXT: .value_type: i64
366 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
367 ; CHECK-NEXT: .offset: 16
368 ; CHECK-NEXT: .size: 8
369 ; CHECK-NEXT: .value_type: i64
370 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
371 ; CHECK-NEXT: .offset: 24
372 ; CHECK-NEXT: .size: 8
373 ; CHECK-NEXT: .value_type: i64
374 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
375 ; CHECK-NEXT: .offset: 32
376 ; CHECK-NEXT: .size: 8
377 ; CHECK-NEXT: .value_type: i8
378 ; CHECK-NEXT: .address_space: global
379 define amdgpu_kernel void @test_sampler(i32 %a)
380 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !18
381 !kernel_arg_base_type !18 !kernel_arg_type_qual !4 {
382 ret void
383 }
385 ; CHECK: .symbol: test_queue.kd
386 ; CHECK: .name: test_queue
387 ; CHECK: .language: OpenCL C
388 ; CHECK: .language_version:
389 ; CHECK-NEXT: - 2
390 ; CHECK-NEXT: - 0
391 ; CHECK: .args:
392 ; CHECK-NEXT: - .type_name: queue_t
393 ; CHECK-NEXT: .value_kind: queue
394 ; CHECK-NEXT: .name: a
395 ; CHECK-NEXT: .offset: 0
396 ; CHECK-NEXT: .size: 8
397 ; CHECK-NEXT: .value_type: struct
398 ; CHECK-NEXT: .address_space: global
399 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
400 ; CHECK-NEXT: .offset: 8
401 ; CHECK-NEXT: .size: 8
402 ; CHECK-NEXT: .value_type: i64
403 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
404 ; CHECK-NEXT: .offset: 16
405 ; CHECK-NEXT: .size: 8
406 ; CHECK-NEXT: .value_type: i64
407 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
408 ; CHECK-NEXT: .offset: 24
409 ; CHECK-NEXT: .size: 8
410 ; CHECK-NEXT: .value_type: i64
411 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
412 ; CHECK-NEXT: .offset: 32
413 ; CHECK-NEXT: .size: 8
414 ; CHECK-NEXT: .value_type: i8
415 ; CHECK-NEXT: .address_space: global
416 define amdgpu_kernel void @test_queue(%opencl.queue_t addrspace(1)* %a)
417 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !19
418 !kernel_arg_base_type !19 !kernel_arg_type_qual !4 {
419 ret void
420 }
422 ; CHECK: .symbol: test_struct.kd
423 ; CHECK: .name: test_struct
424 ; CHECK: .language: OpenCL C
425 ; CHECK: .language_version:
426 ; CHECK-NEXT: - 2
427 ; CHECK-NEXT: - 0
428 ; CHECK: .args:
429 ; CHECK-NEXT: - .type_name: struct A
430 ; CHECK-NEXT: .value_kind: global_buffer
431 ; CHECK-NEXT: .name: a
432 ; CHECK-NEXT: .offset: 0
433 ; CHECK-NEXT: .size: 4
434 ; CHECK-NEXT: .value_type: struct
435 ; CHECK-NEXT: .address_space: private
436 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
437 ; CHECK-NEXT: .offset: 8
438 ; CHECK-NEXT: .size: 8
439 ; CHECK-NEXT: .value_type: i64
440 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
441 ; CHECK-NEXT: .offset: 16
442 ; CHECK-NEXT: .size: 8
443 ; CHECK-NEXT: .value_type: i64
444 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
445 ; CHECK-NEXT: .offset: 24
446 ; CHECK-NEXT: .size: 8
447 ; CHECK-NEXT: .value_type: i64
448 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
449 ; CHECK-NEXT: .offset: 32
450 ; CHECK-NEXT: .size: 8
451 ; CHECK-NEXT: .value_type: i8
452 ; CHECK-NEXT: .address_space: global
453 define amdgpu_kernel void @test_struct(%struct.A addrspace(5)* byval %a)
454 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20
455 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 {
456 ret void
457 }
459 ; CHECK: .symbol: test_i128.kd
460 ; CHECK: .name: test_i128
461 ; CHECK: .language: OpenCL C
462 ; CHECK: .language_version:
463 ; CHECK-NEXT: - 2
464 ; CHECK-NEXT: - 0
465 ; CHECK: .args:
466 ; CHECK-NEXT: - .type_name: i128
467 ; CHECK-NEXT: .value_kind: by_value
468 ; CHECK-NEXT: .offset: 0
469 ; CHECK-NEXT: .size: 16
470 ; CHECK-NEXT: .value_type: struct
471 ; CHECK-NEXT: .name: a
472 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
473 ; CHECK-NEXT: .offset: 16
474 ; CHECK-NEXT: .size: 8
475 ; CHECK-NEXT: .value_type: i64
476 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
477 ; CHECK-NEXT: .offset: 24
478 ; CHECK-NEXT: .size: 8
479 ; CHECK-NEXT: .value_type: i64
480 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
481 ; CHECK-NEXT: .offset: 32
482 ; CHECK-NEXT: .size: 8
483 ; CHECK-NEXT: .value_type: i64
484 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
485 ; CHECK-NEXT: .offset: 40
486 ; CHECK-NEXT: .size: 8
487 ; CHECK-NEXT: .value_type: i8
488 ; CHECK-NEXT: .address_space: global
489 define amdgpu_kernel void @test_i128(i128 %a)
490 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !21
491 !kernel_arg_base_type !21 !kernel_arg_type_qual !4 {
492 ret void
493 }
495 ; CHECK: .symbol: test_multi_arg.kd
496 ; CHECK: .name: test_multi_arg
497 ; CHECK: .language: OpenCL C
498 ; CHECK: .language_version:
499 ; CHECK-NEXT: - 2
500 ; CHECK-NEXT: - 0
501 ; CHECK: .args:
502 ; CHECK-NEXT: - .type_name: int
503 ; CHECK-NEXT: .value_kind: by_value
504 ; CHECK-NEXT: .offset: 0
505 ; CHECK-NEXT: .size: 4
506 ; CHECK-NEXT: .value_type: i32
507 ; CHECK-NEXT: .name: a
508 ; CHECK-NEXT: - .type_name: short2
509 ; CHECK-NEXT: .value_kind: by_value
510 ; CHECK-NEXT: .offset: 4
511 ; CHECK-NEXT: .size: 4
512 ; CHECK-NEXT: .value_type: i16
513 ; CHECK-NEXT: .name: b
514 ; CHECK-NEXT: - .type_name: char3
515 ; CHECK-NEXT: .value_kind: by_value
516 ; CHECK-NEXT: .offset: 8
517 ; CHECK-NEXT: .size: 4
518 ; CHECK-NEXT: .value_type: i8
519 ; CHECK-NEXT: .name: c
520 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
521 ; CHECK-NEXT: .offset: 16
522 ; CHECK-NEXT: .size: 8
523 ; CHECK-NEXT: .value_type: i64
524 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
525 ; CHECK-NEXT: .offset: 24
526 ; CHECK-NEXT: .size: 8
527 ; CHECK-NEXT: .value_type: i64
528 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
529 ; CHECK-NEXT: .offset: 32
530 ; CHECK-NEXT: .size: 8
531 ; CHECK-NEXT: .value_type: i64
532 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
533 ; CHECK-NEXT: .offset: 40
534 ; CHECK-NEXT: .size: 8
535 ; CHECK-NEXT: .value_type: i8
536 ; CHECK-NEXT: .address_space: global
537 define amdgpu_kernel void @test_multi_arg(i32 %a, <2 x i16> %b, <3 x i8> %c)
538 !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !24
539 !kernel_arg_base_type !24 !kernel_arg_type_qual !25 {
540 ret void
541 }
543 ; CHECK: .symbol: test_addr_space.kd
544 ; CHECK: .name: test_addr_space
545 ; CHECK: .language: OpenCL C
546 ; CHECK: .language_version:
547 ; CHECK-NEXT: - 2
548 ; CHECK-NEXT: - 0
549 ; CHECK: .args:
550 ; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
551 ; CHECK-NEXT: .value_kind: global_buffer
552 ; CHECK-NEXT: .name: g
553 ; CHECK-NEXT: .offset: 0
554 ; CHECK-NEXT: .size: 8
555 ; CHECK-NEXT: .value_type: i32
556 ; CHECK-NEXT: .address_space: global
557 ; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
558 ; CHECK-NEXT: .value_kind: global_buffer
559 ; CHECK-NEXT: .name: c
560 ; CHECK-NEXT: .offset: 8
561 ; CHECK-NEXT: .size: 8
562 ; CHECK-NEXT: .value_type: i32
563 ; CHECK-NEXT: .address_space: constant
564 ; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
565 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
566 ; CHECK-NEXT: .name: l
567 ; CHECK-NEXT: .offset: 16
568 ; CHECK-NEXT: .size: 4
569 ; CHECK-NEXT: .value_type: i32
570 ; CHECK-NEXT: .pointee_align: 4
571 ; CHECK-NEXT: .address_space: local
572 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
573 ; CHECK-NEXT: .offset: 24
574 ; CHECK-NEXT: .size: 8
575 ; CHECK-NEXT: .value_type: i64
576 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
577 ; CHECK-NEXT: .offset: 32
578 ; CHECK-NEXT: .size: 8
579 ; CHECK-NEXT: .value_type: i64
580 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
581 ; CHECK-NEXT: .offset: 40
582 ; CHECK-NEXT: .size: 8
583 ; CHECK-NEXT: .value_type: i64
584 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
585 ; CHECK-NEXT: .offset: 48
586 ; CHECK-NEXT: .size: 8
587 ; CHECK-NEXT: .value_type: i8
588 ; CHECK-NEXT: .address_space: global
589 define amdgpu_kernel void @test_addr_space(i32 addrspace(1)* %g,
590 i32 addrspace(4)* %c,
591 i32 addrspace(3)* %l)
592 !kernel_arg_addr_space !50 !kernel_arg_access_qual !23 !kernel_arg_type !51
593 !kernel_arg_base_type !51 !kernel_arg_type_qual !25 {
594 ret void
595 }
597 ; CHECK: .symbol: test_type_qual.kd
598 ; CHECK: .name: test_type_qual
599 ; CHECK: .language: OpenCL C
600 ; CHECK: .language_version:
601 ; CHECK-NEXT: - 2
602 ; CHECK-NEXT: - 0
603 ; CHECK: .args:
604 ; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
605 ; CHECK-NEXT: .value_kind: global_buffer
606 ; CHECK-NEXT: .name: a
607 ; CHECK-NEXT: .offset: 0
608 ; CHECK-NEXT: .size: 8
609 ; CHECK-NEXT: .is_volatile: true
610 ; CHECK-NEXT: .value_type: i32
611 ; CHECK-NEXT: .address_space: global
612 ; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
613 ; CHECK-NEXT: .value_kind: global_buffer
614 ; CHECK-NEXT: .name: b
615 ; CHECK-NEXT: .is_const: true
616 ; CHECK-NEXT: .offset: 8
617 ; CHECK-NEXT: .size: 8
618 ; CHECK-NEXT: .is_restrict: true
619 ; CHECK-NEXT: .value_type: i32
620 ; CHECK-NEXT: .address_space: global
621 ; CHECK-NEXT: - .type_name: 'int addrspace(5)*'
622 ; CHECK-NEXT: .value_kind: pipe
623 ; CHECK-NEXT: .name: c
624 ; CHECK-NEXT: .offset: 16
625 ; CHECK-NEXT: .is_pipe: true
626 ; CHECK-NEXT: .size: 8
627 ; CHECK-NEXT: .value_type: struct
628 ; CHECK-NEXT: .address_space: global
629 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
630 ; CHECK-NEXT: .offset: 24
631 ; CHECK-NEXT: .size: 8
632 ; CHECK-NEXT: .value_type: i64
633 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
634 ; CHECK-NEXT: .offset: 32
635 ; CHECK-NEXT: .size: 8
636 ; CHECK-NEXT: .value_type: i64
637 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
638 ; CHECK-NEXT: .offset: 40
639 ; CHECK-NEXT: .size: 8
640 ; CHECK-NEXT: .value_type: i64
641 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
642 ; CHECK-NEXT: .offset: 48
643 ; CHECK-NEXT: .size: 8
644 ; CHECK-NEXT: .value_type: i8
645 ; CHECK-NEXT: .address_space: global
646 define amdgpu_kernel void @test_type_qual(i32 addrspace(1)* %a,
647 i32 addrspace(1)* %b,
648 %opencl.pipe_t addrspace(1)* %c)
649 !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !51
650 !kernel_arg_base_type !51 !kernel_arg_type_qual !70 {
651 ret void
652 }
654 ; CHECK: .symbol: test_access_qual.kd
655 ; CHECK: .name: test_access_qual
656 ; CHECK: .language: OpenCL C
657 ; CHECK: .language_version:
658 ; CHECK-NEXT: - 2
659 ; CHECK-NEXT: - 0
660 ; CHECK: .args:
661 ; CHECK-NEXT: - .type_name: image1d_t
662 ; CHECK-NEXT: .value_kind: image
663 ; CHECK-NEXT: .name: ro
664 ; CHECK-NEXT: .access: read_only
665 ; CHECK-NEXT: .offset: 0
666 ; CHECK-NEXT: .size: 8
667 ; CHECK-NEXT: .value_type: struct
668 ; CHECK-NEXT: .address_space: global
669 ; CHECK-NEXT: - .type_name: image2d_t
670 ; CHECK-NEXT: .value_kind: image
671 ; CHECK-NEXT: .name: wo
672 ; CHECK-NEXT: .access: write_only
673 ; CHECK-NEXT: .offset: 8
674 ; CHECK-NEXT: .size: 8
675 ; CHECK-NEXT: .value_type: struct
676 ; CHECK-NEXT: .address_space: global
677 ; CHECK-NEXT: - .type_name: image3d_t
678 ; CHECK-NEXT: .value_kind: image
679 ; CHECK-NEXT: .name: rw
680 ; CHECK-NEXT: .access: read_write
681 ; CHECK-NEXT: .offset: 16
682 ; CHECK-NEXT: .size: 8
683 ; CHECK-NEXT: .value_type: struct
684 ; CHECK-NEXT: .address_space: global
685 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
686 ; CHECK-NEXT: .offset: 24
687 ; CHECK-NEXT: .size: 8
688 ; CHECK-NEXT: .value_type: i64
689 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
690 ; CHECK-NEXT: .offset: 32
691 ; CHECK-NEXT: .size: 8
692 ; CHECK-NEXT: .value_type: i64
693 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
694 ; CHECK-NEXT: .offset: 40
695 ; CHECK-NEXT: .size: 8
696 ; CHECK-NEXT: .value_type: i64
697 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
698 ; CHECK-NEXT: .offset: 48
699 ; CHECK-NEXT: .size: 8
700 ; CHECK-NEXT: .value_type: i8
701 ; CHECK-NEXT: .address_space: global
702 define amdgpu_kernel void @test_access_qual(%opencl.image1d_t addrspace(1)* %ro,
703 %opencl.image2d_t addrspace(1)* %wo,
704 %opencl.image3d_t addrspace(1)* %rw)
705 !kernel_arg_addr_space !60 !kernel_arg_access_qual !61 !kernel_arg_type !62
706 !kernel_arg_base_type !62 !kernel_arg_type_qual !25 {
707 ret void
708 }
710 ; CHECK: .symbol: test_vec_type_hint_half.kd
711 ; CHECK: .name: test_vec_type_hint_half
712 ; CHECK: .language: OpenCL C
713 ; CHECK: .language_version:
714 ; CHECK-NEXT: - 2
715 ; CHECK-NEXT: - 0
716 ; CHECK: .args:
717 ; CHECK-NEXT: - .type_name: int
718 ; CHECK-NEXT: .value_kind: by_value
719 ; CHECK-NEXT: .offset: 0
720 ; CHECK-NEXT: .size: 4
721 ; CHECK-NEXT: .value_type: i32
722 ; CHECK-NEXT: .name: a
723 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
724 ; CHECK-NEXT: .offset: 8
725 ; CHECK-NEXT: .size: 8
726 ; CHECK-NEXT: .value_type: i64
727 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
728 ; CHECK-NEXT: .offset: 16
729 ; CHECK-NEXT: .size: 8
730 ; CHECK-NEXT: .value_type: i64
731 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
732 ; CHECK-NEXT: .offset: 24
733 ; CHECK-NEXT: .size: 8
734 ; CHECK-NEXT: .value_type: i64
735 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
736 ; CHECK-NEXT: .offset: 32
737 ; CHECK-NEXT: .size: 8
738 ; CHECK-NEXT: .value_type: i8
739 ; CHECK-NEXT: .address_space: global
740 ; CHECK: .vec_type_hint: half
741 define amdgpu_kernel void @test_vec_type_hint_half(i32 %a)
742 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
743 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !26 {
744 ret void
745 }
747 ; CHECK: .symbol: test_vec_type_hint_float.kd
748 ; CHECK: .name: test_vec_type_hint_float
749 ; CHECK: .language: OpenCL C
750 ; CHECK: .language_version:
751 ; CHECK-NEXT: - 2
752 ; CHECK-NEXT: - 0
753 ; CHECK: .args:
754 ; CHECK-NEXT: - .type_name: int
755 ; CHECK-NEXT: .value_kind: by_value
756 ; CHECK-NEXT: .offset: 0
757 ; CHECK-NEXT: .size: 4
758 ; CHECK-NEXT: .value_type: i32
759 ; CHECK-NEXT: .name: a
760 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
761 ; CHECK-NEXT: .offset: 8
762 ; CHECK-NEXT: .size: 8
763 ; CHECK-NEXT: .value_type: i64
764 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
765 ; CHECK-NEXT: .offset: 16
766 ; CHECK-NEXT: .size: 8
767 ; CHECK-NEXT: .value_type: i64
768 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
769 ; CHECK-NEXT: .offset: 24
770 ; CHECK-NEXT: .size: 8
771 ; CHECK-NEXT: .value_type: i64
772 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
773 ; CHECK-NEXT: .offset: 32
774 ; CHECK-NEXT: .size: 8
775 ; CHECK-NEXT: .value_type: i8
776 ; CHECK-NEXT: .address_space: global
777 ; CHECK: .vec_type_hint: float
778 define amdgpu_kernel void @test_vec_type_hint_float(i32 %a)
779 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
780 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !27 {
781 ret void
782 }
784 ; CHECK: .symbol: test_vec_type_hint_double.kd
785 ; CHECK: .name: test_vec_type_hint_double
786 ; CHECK: .language: OpenCL C
787 ; CHECK: .language_version:
788 ; CHECK-NEXT: - 2
789 ; CHECK-NEXT: - 0
790 ; CHECK: .args:
791 ; CHECK-NEXT: - .type_name: int
792 ; CHECK-NEXT: .value_kind: by_value
793 ; CHECK-NEXT: .offset: 0
794 ; CHECK-NEXT: .size: 4
795 ; CHECK-NEXT: .value_type: i32
796 ; CHECK-NEXT: .name: a
797 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
798 ; CHECK-NEXT: .offset: 8
799 ; CHECK-NEXT: .size: 8
800 ; CHECK-NEXT: .value_type: i64
801 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
802 ; CHECK-NEXT: .offset: 16
803 ; CHECK-NEXT: .size: 8
804 ; CHECK-NEXT: .value_type: i64
805 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
806 ; CHECK-NEXT: .offset: 24
807 ; CHECK-NEXT: .size: 8
808 ; CHECK-NEXT: .value_type: i64
809 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
810 ; CHECK-NEXT: .offset: 32
811 ; CHECK-NEXT: .size: 8
812 ; CHECK-NEXT: .value_type: i8
813 ; CHECK-NEXT: .address_space: global
814 ; CHECK: .vec_type_hint: double
815 define amdgpu_kernel void @test_vec_type_hint_double(i32 %a)
816 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
817 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !28 {
818 ret void
819 }
821 ; CHECK: .symbol: test_vec_type_hint_char.kd
822 ; CHECK: .name: test_vec_type_hint_char
823 ; CHECK: .language: OpenCL C
824 ; CHECK: .language_version:
825 ; CHECK-NEXT: - 2
826 ; CHECK-NEXT: - 0
827 ; CHECK: .args:
828 ; CHECK-NEXT: - .type_name: int
829 ; CHECK-NEXT: .value_kind: by_value
830 ; CHECK-NEXT: .offset: 0
831 ; CHECK-NEXT: .size: 4
832 ; CHECK-NEXT: .value_type: i32
833 ; CHECK-NEXT: .name: a
834 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
835 ; CHECK-NEXT: .offset: 8
836 ; CHECK-NEXT: .size: 8
837 ; CHECK-NEXT: .value_type: i64
838 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
839 ; CHECK-NEXT: .offset: 16
840 ; CHECK-NEXT: .size: 8
841 ; CHECK-NEXT: .value_type: i64
842 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
843 ; CHECK-NEXT: .offset: 24
844 ; CHECK-NEXT: .size: 8
845 ; CHECK-NEXT: .value_type: i64
846 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
847 ; CHECK-NEXT: .offset: 32
848 ; CHECK-NEXT: .size: 8
849 ; CHECK-NEXT: .value_type: i8
850 ; CHECK-NEXT: .address_space: global
851 ; CHECK: .vec_type_hint: char
852 define amdgpu_kernel void @test_vec_type_hint_char(i32 %a)
853 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
854 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !29 {
855 ret void
856 }
858 ; CHECK: .symbol: test_vec_type_hint_short.kd
859 ; CHECK: .name: test_vec_type_hint_short
860 ; CHECK: .language: OpenCL C
861 ; CHECK: .language_version:
862 ; CHECK-NEXT: - 2
863 ; CHECK-NEXT: - 0
864 ; CHECK: .args:
865 ; CHECK-NEXT: - .type_name: int
866 ; CHECK-NEXT: .value_kind: by_value
867 ; CHECK-NEXT: .offset: 0
868 ; CHECK-NEXT: .size: 4
869 ; CHECK-NEXT: .value_type: i32
870 ; CHECK-NEXT: .name: a
871 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
872 ; CHECK-NEXT: .offset: 8
873 ; CHECK-NEXT: .size: 8
874 ; CHECK-NEXT: .value_type: i64
875 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
876 ; CHECK-NEXT: .offset: 16
877 ; CHECK-NEXT: .size: 8
878 ; CHECK-NEXT: .value_type: i64
879 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
880 ; CHECK-NEXT: .offset: 24
881 ; CHECK-NEXT: .size: 8
882 ; CHECK-NEXT: .value_type: i64
883 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
884 ; CHECK-NEXT: .offset: 32
885 ; CHECK-NEXT: .size: 8
886 ; CHECK-NEXT: .value_type: i8
887 ; CHECK-NEXT: .address_space: global
888 ; CHECK: .vec_type_hint: short
889 define amdgpu_kernel void @test_vec_type_hint_short(i32 %a)
890 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
891 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !30 {
892 ret void
893 }
895 ; CHECK: .symbol: test_vec_type_hint_long.kd
896 ; CHECK: .name: test_vec_type_hint_long
897 ; CHECK: .language: OpenCL C
898 ; CHECK: .language_version:
899 ; CHECK-NEXT: - 2
900 ; CHECK-NEXT: - 0
901 ; CHECK: .args:
902 ; CHECK-NEXT: - .type_name: int
903 ; CHECK-NEXT: .value_kind: by_value
904 ; CHECK-NEXT: .offset: 0
905 ; CHECK-NEXT: .size: 4
906 ; CHECK-NEXT: .value_type: i32
907 ; CHECK-NEXT: .name: a
908 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
909 ; CHECK-NEXT: .offset: 8
910 ; CHECK-NEXT: .size: 8
911 ; CHECK-NEXT: .value_type: i64
912 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
913 ; CHECK-NEXT: .offset: 16
914 ; CHECK-NEXT: .size: 8
915 ; CHECK-NEXT: .value_type: i64
916 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
917 ; CHECK-NEXT: .offset: 24
918 ; CHECK-NEXT: .size: 8
919 ; CHECK-NEXT: .value_type: i64
920 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
921 ; CHECK-NEXT: .offset: 32
922 ; CHECK-NEXT: .size: 8
923 ; CHECK-NEXT: .value_type: i8
924 ; CHECK-NEXT: .address_space: global
925 ; CHECK: .vec_type_hint: long
926 define amdgpu_kernel void @test_vec_type_hint_long(i32 %a)
927 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
928 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !31 {
929 ret void
930 }
932 ; CHECK: .symbol: test_vec_type_hint_unknown.kd
933 ; CHECK: .name: test_vec_type_hint_unknown
934 ; CHECK: .language: OpenCL C
935 ; CHECK: .language_version:
936 ; CHECK-NEXT: - 2
937 ; CHECK-NEXT: - 0
938 ; CHECK: .args:
939 ; CHECK-NEXT: - .type_name: int
940 ; CHECK-NEXT: .value_kind: by_value
941 ; CHECK-NEXT: .offset: 0
942 ; CHECK-NEXT: .size: 4
943 ; CHECK-NEXT: .value_type: i32
944 ; CHECK-NEXT: .name: a
945 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
946 ; CHECK-NEXT: .offset: 8
947 ; CHECK-NEXT: .size: 8
948 ; CHECK-NEXT: .value_type: i64
949 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
950 ; CHECK-NEXT: .offset: 16
951 ; CHECK-NEXT: .size: 8
952 ; CHECK-NEXT: .value_type: i64
953 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
954 ; CHECK-NEXT: .offset: 24
955 ; CHECK-NEXT: .size: 8
956 ; CHECK-NEXT: .value_type: i64
957 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
958 ; CHECK-NEXT: .offset: 32
959 ; CHECK-NEXT: .size: 8
960 ; CHECK-NEXT: .value_type: i8
961 ; CHECK-NEXT: .address_space: global
962 ; CHECK: .vec_type_hint: unknown
963 define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %a)
964 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
965 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !32 {
966 ret void
967 }
969 ; CHECK: .reqd_workgroup_size:
970 ; CHECK-NEXT: - 1
971 ; CHECK-NEXT: - 2
972 ; CHECK-NEXT: - 4
973 ; CHECK: .symbol: test_reqd_wgs_vec_type_hint.kd
974 ; CHECK: .name: test_reqd_wgs_vec_type_hint
975 ; CHECK: .language: OpenCL C
976 ; CHECK: .language_version:
977 ; CHECK-NEXT: - 2
978 ; CHECK-NEXT: - 0
979 ; CHECK: .args:
980 ; CHECK-NEXT: - .type_name: int
981 ; CHECK-NEXT: .value_kind: by_value
982 ; CHECK-NEXT: .offset: 0
983 ; CHECK-NEXT: .size: 4
984 ; CHECK-NEXT: .value_type: i32
985 ; CHECK-NEXT: .name: a
986 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
987 ; CHECK-NEXT: .offset: 8
988 ; CHECK-NEXT: .size: 8
989 ; CHECK-NEXT: .value_type: i64
990 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
991 ; CHECK-NEXT: .offset: 16
992 ; CHECK-NEXT: .size: 8
993 ; CHECK-NEXT: .value_type: i64
994 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
995 ; CHECK-NEXT: .offset: 24
996 ; CHECK-NEXT: .size: 8
997 ; CHECK-NEXT: .value_type: i64
998 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
999 ; CHECK-NEXT: .offset: 32
1000 ; CHECK-NEXT: .size: 8
1001 ; CHECK-NEXT: .value_type: i8
1002 ; CHECK-NEXT: .address_space: global
1003 ; CHECK: .vec_type_hint: int
1004 define amdgpu_kernel void @test_reqd_wgs_vec_type_hint(i32 %a)
1005 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
1006 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !5
1007 !reqd_work_group_size !6 {
1008 ret void
1009 }
1011 ; CHECK: .symbol: test_wgs_hint_vec_type_hint.kd
1012 ; CHECK: .workgroup_size_hint:
1013 ; CHECK-NEXT: - 8
1014 ; CHECK-NEXT: - 16
1015 ; CHECK-NEXT: - 32
1016 ; CHECK: .name: test_wgs_hint_vec_type_hint
1017 ; CHECK: .language: OpenCL C
1018 ; CHECK: .language_version:
1019 ; CHECK-NEXT: - 2
1020 ; CHECK-NEXT: - 0
1021 ; CHECK: .args:
1022 ; CHECK-NEXT: - .type_name: int
1023 ; CHECK-NEXT: .value_kind: by_value
1024 ; CHECK-NEXT: .offset: 0
1025 ; CHECK-NEXT: .size: 4
1026 ; CHECK-NEXT: .value_type: i32
1027 ; CHECK-NEXT: .name: a
1028 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
1029 ; CHECK-NEXT: .offset: 8
1030 ; CHECK-NEXT: .size: 8
1031 ; CHECK-NEXT: .value_type: i64
1032 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
1033 ; CHECK-NEXT: .offset: 16
1034 ; CHECK-NEXT: .size: 8
1035 ; CHECK-NEXT: .value_type: i64
1036 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
1037 ; CHECK-NEXT: .offset: 24
1038 ; CHECK-NEXT: .size: 8
1039 ; CHECK-NEXT: .value_type: i64
1040 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
1041 ; CHECK-NEXT: .offset: 32
1042 ; CHECK-NEXT: .size: 8
1043 ; CHECK-NEXT: .value_type: i8
1044 ; CHECK-NEXT: .address_space: global
1045 ; CHECK: .vec_type_hint: uint4
1046 define amdgpu_kernel void @test_wgs_hint_vec_type_hint(i32 %a)
1047 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
1048 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !7
1049 !work_group_size_hint !8 {
1050 ret void
1051 }
1053 ; CHECK: .symbol: test_arg_ptr_to_ptr.kd
1054 ; CHECK: .name: test_arg_ptr_to_ptr
1055 ; CHECK: .language: OpenCL C
1056 ; CHECK: .language_version:
1057 ; CHECK-NEXT: - 2
1058 ; CHECK-NEXT: - 0
1059 ; CHECK: .args:
1060 ; CHECK-NEXT: - .type_name: 'int addrspace(5)* addrspace(5)*'
1061 ; CHECK-NEXT: .value_kind: global_buffer
1062 ; CHECK-NEXT: .name: a
1063 ; CHECK-NEXT: .offset: 0
1064 ; CHECK-NEXT: .size: 8
1065 ; CHECK-NEXT: .value_type: i32
1066 ; CHECK-NEXT: .address_space: global
1067 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
1068 ; CHECK-NEXT: .offset: 8
1069 ; CHECK-NEXT: .size: 8
1070 ; CHECK-NEXT: .value_type: i64
1071 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
1072 ; CHECK-NEXT: .offset: 16
1073 ; CHECK-NEXT: .size: 8
1074 ; CHECK-NEXT: .value_type: i64
1075 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
1076 ; CHECK-NEXT: .offset: 24
1077 ; CHECK-NEXT: .size: 8
1078 ; CHECK-NEXT: .value_type: i64
1079 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
1080 ; CHECK-NEXT: .offset: 32
1081 ; CHECK-NEXT: .size: 8
1082 ; CHECK-NEXT: .value_type: i8
1083 ; CHECK-NEXT: .address_space: global
1084 define amdgpu_kernel void @test_arg_ptr_to_ptr(i32 addrspace(5)* addrspace(1)* %a)
1085 !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !80
1086 !kernel_arg_base_type !80 !kernel_arg_type_qual !4 {
1087 ret void
1088 }
1090 ; CHECK: .symbol: test_arg_struct_contains_ptr.kd
1091 ; CHECK: .name: test_arg_struct_contains_ptr
1092 ; CHECK: .language: OpenCL C
1093 ; CHECK: .language_version:
1094 ; CHECK-NEXT: - 2
1095 ; CHECK-NEXT: - 0
1096 ; CHECK: .args:
1097 ; CHECK-NEXT: - .type_name: struct B
1098 ; CHECK-NEXT: .value_kind: global_buffer
1099 ; CHECK-NEXT: .name: a
1100 ; CHECK-NEXT: .offset: 0
1101 ; CHECK-NEXT: .size: 4
1102 ; CHECK-NEXT: .value_type: struct
1103 ; CHECK-NEXT: .address_space: private
1104 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
1105 ; CHECK-NEXT: .offset: 8
1106 ; CHECK-NEXT: .size: 8
1107 ; CHECK-NEXT: .value_type: i64
1108 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
1109 ; CHECK-NEXT: .offset: 16
1110 ; CHECK-NEXT: .size: 8
1111 ; CHECK-NEXT: .value_type: i64
1112 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
1113 ; CHECK-NEXT: .offset: 24
1114 ; CHECK-NEXT: .size: 8
1115 ; CHECK-NEXT: .value_type: i64
1116 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
1117 ; CHECK-NEXT: .offset: 32
1118 ; CHECK-NEXT: .size: 8
1119 ; CHECK-NEXT: .value_type: i8
1120 ; CHECK-NEXT: .address_space: global
1121 define amdgpu_kernel void @test_arg_struct_contains_ptr(%struct.B addrspace(5)* byval %a)
1122 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !82
1123 !kernel_arg_base_type !82 !kernel_arg_type_qual !4 {
1124 ret void
1125 }
1127 ; CHECK: .symbol: test_arg_vector_of_ptr.kd
1128 ; CHECK: .name: test_arg_vector_of_ptr
1129 ; CHECK: .language: OpenCL C
1130 ; CHECK: .language_version:
1131 ; CHECK-NEXT: - 2
1132 ; CHECK-NEXT: - 0
1133 ; CHECK: .args:
1134 ; CHECK-NEXT: - .type_name: 'global int addrspace(5)* __attribute__((ext_vector_type(2)))'
1135 ; CHECK-NEXT: .value_kind: by_value
1136 ; CHECK-NEXT: .offset: 0
1137 ; CHECK-NEXT: .size: 16
1138 ; CHECK-NEXT: .value_type: i32
1139 ; CHECK-NEXT: .name: a
1140 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
1141 ; CHECK-NEXT: .offset: 16
1142 ; CHECK-NEXT: .size: 8
1143 ; CHECK-NEXT: .value_type: i64
1144 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
1145 ; CHECK-NEXT: .offset: 24
1146 ; CHECK-NEXT: .size: 8
1147 ; CHECK-NEXT: .value_type: i64
1148 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
1149 ; CHECK-NEXT: .offset: 32
1150 ; CHECK-NEXT: .size: 8
1151 ; CHECK-NEXT: .value_type: i64
1152 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
1153 ; CHECK-NEXT: .offset: 40
1154 ; CHECK-NEXT: .size: 8
1155 ; CHECK-NEXT: .value_type: i8
1156 ; CHECK-NEXT: .address_space: global
1157 define amdgpu_kernel void @test_arg_vector_of_ptr(<2 x i32 addrspace(1)*> %a)
1158 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !83
1159 !kernel_arg_base_type !83 !kernel_arg_type_qual !4 {
1160 ret void
1161 }
1163 ; CHECK: .symbol: test_arg_unknown_builtin_type.kd
1164 ; CHECK: .name: test_arg_unknown_builtin_type
1165 ; CHECK: .language: OpenCL C
1166 ; CHECK: .language_version:
1167 ; CHECK-NEXT: - 2
1168 ; CHECK-NEXT: - 0
1169 ; CHECK: .args:
1170 ; CHECK-NEXT: - .type_name: clk_event_t
1171 ; CHECK-NEXT: .value_kind: global_buffer
1172 ; CHECK-NEXT: .name: a
1173 ; CHECK-NEXT: .offset: 0
1174 ; CHECK-NEXT: .size: 8
1175 ; CHECK-NEXT: .value_type: struct
1176 ; CHECK-NEXT: .address_space: global
1177 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
1178 ; CHECK-NEXT: .offset: 8
1179 ; CHECK-NEXT: .size: 8
1180 ; CHECK-NEXT: .value_type: i64
1181 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
1182 ; CHECK-NEXT: .offset: 16
1183 ; CHECK-NEXT: .size: 8
1184 ; CHECK-NEXT: .value_type: i64
1185 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
1186 ; CHECK-NEXT: .offset: 24
1187 ; CHECK-NEXT: .size: 8
1188 ; CHECK-NEXT: .value_type: i64
1189 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
1190 ; CHECK-NEXT: .offset: 32
1191 ; CHECK-NEXT: .size: 8
1192 ; CHECK-NEXT: .value_type: i8
1193 ; CHECK-NEXT: .address_space: global
1194 define amdgpu_kernel void @test_arg_unknown_builtin_type(
1195 %opencl.clk_event_t addrspace(1)* %a)
1196 !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !84
1197 !kernel_arg_base_type !84 !kernel_arg_type_qual !4 {
1198 ret void
1199 }
1201 ; CHECK: .symbol: test_pointee_align.kd
1202 ; CHECK: .name: test_pointee_align
1203 ; CHECK: .language: OpenCL C
1204 ; CHECK: .language_version:
1205 ; CHECK-NEXT: - 2
1206 ; CHECK-NEXT: - 0
1207 ; CHECK: .args:
1208 ; CHECK-NEXT: - .type_name: 'long addrspace(5)*'
1209 ; CHECK-NEXT: .value_kind: global_buffer
1210 ; CHECK-NEXT: .name: a
1211 ; CHECK-NEXT: .offset: 0
1212 ; CHECK-NEXT: .size: 8
1213 ; CHECK-NEXT: .value_type: i64
1214 ; CHECK-NEXT: .address_space: global
1215 ; CHECK-NEXT: - .type_name: 'char addrspace(5)*'
1216 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1217 ; CHECK-NEXT: .name: b
1218 ; CHECK-NEXT: .offset: 8
1219 ; CHECK-NEXT: .size: 4
1220 ; CHECK-NEXT: .value_type: i8
1221 ; CHECK-NEXT: .pointee_align: 1
1222 ; CHECK-NEXT: .address_space: local
1223 ; CHECK-NEXT: - .type_name: 'char2 addrspace(5)*'
1224 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1225 ; CHECK-NEXT: .name: c
1226 ; CHECK-NEXT: .offset: 12
1227 ; CHECK-NEXT: .size: 4
1228 ; CHECK-NEXT: .value_type: i8
1229 ; CHECK-NEXT: .pointee_align: 2
1230 ; CHECK-NEXT: .address_space: local
1231 ; CHECK-NEXT: - .type_name: 'char3 addrspace(5)*'
1232 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1233 ; CHECK-NEXT: .name: d
1234 ; CHECK-NEXT: .offset: 16
1235 ; CHECK-NEXT: .size: 4
1236 ; CHECK-NEXT: .value_type: i8
1237 ; CHECK-NEXT: .pointee_align: 4
1238 ; CHECK-NEXT: .address_space: local
1239 ; CHECK-NEXT: - .type_name: 'char4 addrspace(5)*'
1240 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1241 ; CHECK-NEXT: .name: e
1242 ; CHECK-NEXT: .offset: 20
1243 ; CHECK-NEXT: .size: 4
1244 ; CHECK-NEXT: .value_type: i8
1245 ; CHECK-NEXT: .pointee_align: 4
1246 ; CHECK-NEXT: .address_space: local
1247 ; CHECK-NEXT: - .type_name: 'char8 addrspace(5)*'
1248 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1249 ; CHECK-NEXT: .name: f
1250 ; CHECK-NEXT: .offset: 24
1251 ; CHECK-NEXT: .size: 4
1252 ; CHECK-NEXT: .value_type: i8
1253 ; CHECK-NEXT: .pointee_align: 8
1254 ; CHECK-NEXT: .address_space: local
1255 ; CHECK-NEXT: - .type_name: 'char16 addrspace(5)*'
1256 ; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1257 ; CHECK-NEXT: .name: g
1258 ; CHECK-NEXT: .offset: 28
1259 ; CHECK-NEXT: .size: 4
1260 ; CHECK-NEXT: .value_type: i8
1261 ; CHECK-NEXT: .pointee_align: 16
1262 ; CHECK-NEXT: .address_space: local
1263 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
1264 ; CHECK-NEXT: .offset: 32
1265 ; CHECK-NEXT: .size: 8
1266 ; CHECK-NEXT: .value_type: i64
1267 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
1268 ; CHECK-NEXT: .offset: 40
1269 ; CHECK-NEXT: .size: 8
1270 ; CHECK-NEXT: .value_type: i64
1271 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
1272 ; CHECK-NEXT: .offset: 48
1273 ; CHECK-NEXT: .size: 8
1274 ; CHECK-NEXT: .value_type: i64
1275 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
1276 ; CHECK-NEXT: .offset: 56
1277 ; CHECK-NEXT: .size: 8
1278 ; CHECK-NEXT: .value_type: i8
1279 ; CHECK-NEXT: .address_space: global
1280 define amdgpu_kernel void @test_pointee_align(i64 addrspace(1)* %a,
1281 i8 addrspace(3)* %b,
1282 <2 x i8> addrspace(3)* %c,
1283 <3 x i8> addrspace(3)* %d,
1284 <4 x i8> addrspace(3)* %e,
1285 <8 x i8> addrspace(3)* %f,
1286 <16 x i8> addrspace(3)* %g)
1287 !kernel_arg_addr_space !91 !kernel_arg_access_qual !92 !kernel_arg_type !93
1288 !kernel_arg_base_type !93 !kernel_arg_type_qual !94 {
1289 ret void
1290 }
1292 ; CHECK: .symbol: __test_block_invoke_kernel.kd
1293 ; CHECK: .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle
1294 ; CHECK: .name: __test_block_invoke_kernel
1295 ; CHECK: .language: OpenCL C
1296 ; CHECK: .language_version:
1297 ; CHECK-NEXT: - 2
1298 ; CHECK-NEXT: - 0
1299 ; CHECK: .args:
1300 ; CHECK-NEXT: - .type_name: __block_literal
1301 ; CHECK-NEXT: .value_kind: by_value
1302 ; CHECK-NEXT: .offset: 0
1303 ; CHECK-NEXT: .size: 25
1304 ; CHECK-NEXT: .value_type: struct
1305 ; CHECK-NEXT: .name: arg
1306 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
1307 ; CHECK-NEXT: .offset: 32
1308 ; CHECK-NEXT: .size: 8
1309 ; CHECK-NEXT: .value_type: i64
1310 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
1311 ; CHECK-NEXT: .offset: 40
1312 ; CHECK-NEXT: .size: 8
1313 ; CHECK-NEXT: .value_type: i64
1314 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
1315 ; CHECK-NEXT: .offset: 48
1316 ; CHECK-NEXT: .size: 8
1317 ; CHECK-NEXT: .value_type: i64
1318 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
1319 ; CHECK-NEXT: .offset: 56
1320 ; CHECK-NEXT: .size: 8
1321 ; CHECK-NEXT: .value_type: i8
1322 ; CHECK-NEXT: .address_space: global
1323 define amdgpu_kernel void @__test_block_invoke_kernel(
1324 <{ i32, i32, i8*, i8 addrspace(1)*, i8 }> %arg) #0
1325 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !110
1326 !kernel_arg_base_type !110 !kernel_arg_type_qual !4 {
1327 ret void
1328 }
1330 ; CHECK: .symbol: test_enqueue_kernel_caller.kd
1331 ; CHECK: .name: test_enqueue_kernel_caller
1332 ; CHECK: .language: OpenCL C
1333 ; CHECK: .language_version:
1334 ; CHECK-NEXT: - 2
1335 ; CHECK-NEXT: - 0
1336 ; CHECK: .args:
1337 ; CHECK-NEXT: - .type_name: char
1338 ; CHECK-NEXT: .value_kind: by_value
1339 ; CHECK-NEXT: .offset: 0
1340 ; CHECK-NEXT: .size: 1
1341 ; CHECK-NEXT: .value_type: i8
1342 ; CHECK-NEXT: .name: a
1343 ; CHECK-NEXT: - .value_kind: hidden_global_offset_x
1344 ; CHECK-NEXT: .offset: 8
1345 ; CHECK-NEXT: .size: 8
1346 ; CHECK-NEXT: .value_type: i64
1347 ; CHECK-NEXT: - .value_kind: hidden_global_offset_y
1348 ; CHECK-NEXT: .offset: 16
1349 ; CHECK-NEXT: .size: 8
1350 ; CHECK-NEXT: .value_type: i64
1351 ; CHECK-NEXT: - .value_kind: hidden_global_offset_z
1352 ; CHECK-NEXT: .offset: 24
1353 ; CHECK-NEXT: .size: 8
1354 ; CHECK-NEXT: .value_type: i64
1355 ; CHECK-NEXT: - .value_kind: hidden_printf_buffer
1356 ; CHECK-NEXT: .offset: 32
1357 ; CHECK-NEXT: .size: 8
1358 ; CHECK-NEXT: .value_type: i8
1359 ; CHECK-NEXT: .address_space: global
1360 ; CHECK-NEXT: - .value_kind: hidden_default_queue
1361 ; CHECK-NEXT: .offset: 40
1362 ; CHECK-NEXT: .size: 8
1363 ; CHECK-NEXT: .value_type: i8
1364 ; CHECK-NEXT: .address_space: global
1365 ; CHECK-NEXT: - .value_kind: hidden_completion_action
1366 ; CHECK-NEXT: .offset: 48
1367 ; CHECK-NEXT: .size: 8
1368 ; CHECK-NEXT: .value_type: i8
1369 ; CHECK-NEXT: .address_space: global
1370 define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #1
1371 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
1372 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
1373 ret void
1374 }
1376 ; CHECK: .symbol: unknown_addrspace_kernarg.kd
1377 ; CHECK: .name: unknown_addrspace_kernarg
1378 ; CHECK: .args:
1379 ; CHECK-NEXT: .value_kind: global_buffer
1380 ; CHECK-NEXT: .offset: 0
1381 ; CHECK-NEXT: .size: 8
1382 ; CHECK-NEXT: .value_type: i32
1383 ; CHECK-NEXT: .name: ptr
1384 define amdgpu_kernel void @unknown_addrspace_kernarg(i32 addrspace(12345)* %ptr) #0 {
1385 ret void
1386 }
1388 ; CHECK: amdhsa.version:
1389 ; CHECK-NEXT: - 1
1390 ; CHECK-NEXT: - 0
1391 ; CHECK: amdhsa.printf:
1392 ; CHECK-NEXT: - '1:1:4:%d\n'
1393 ; CHECK-NEXT: - '2:1:8:%g\n'
1395 attributes #0 = { "runtime-handle"="__test_block_invoke_kernel_runtime_handle" }
1396 attributes #1 = { "calls-enqueue-kernel" }
1398 !llvm.printf.fmts = !{!100, !101}
1400 !1 = !{i32 0}
1401 !2 = !{!"none"}
1402 !3 = !{!"int"}
1403 !4 = !{!""}
1404 !5 = !{i32 undef, i32 1}
1405 !6 = !{i32 1, i32 2, i32 4}
1406 !7 = !{<4 x i32> undef, i32 0}
1407 !8 = !{i32 8, i32 16, i32 32}
1408 !9 = !{!"char"}
1409 !10 = !{!"ushort2"}
1410 !11 = !{!"int3"}
1411 !12 = !{!"ulong4"}
1412 !13 = !{!"half8"}
1413 !14 = !{!"float16"}
1414 !15 = !{!"double16"}
1415 !16 = !{!"int addrspace(5)*"}
1416 !17 = !{!"image2d_t"}
1417 !18 = !{!"sampler_t"}
1418 !19 = !{!"queue_t"}
1419 !20 = !{!"struct A"}
1420 !21 = !{!"i128"}
1421 !22 = !{i32 0, i32 0, i32 0}
1422 !23 = !{!"none", !"none", !"none"}
1423 !24 = !{!"int", !"short2", !"char3"}
1424 !25 = !{!"", !"", !""}
1425 !26 = !{half undef, i32 1}
1426 !27 = !{float undef, i32 1}
1427 !28 = !{double undef, i32 1}
1428 !29 = !{i8 undef, i32 1}
1429 !30 = !{i16 undef, i32 1}
1430 !31 = !{i64 undef, i32 1}
1431 !32 = !{i32 addrspace(5)*undef, i32 1}
1432 !50 = !{i32 1, i32 2, i32 3}
1433 !51 = !{!"int addrspace(5)*", !"int addrspace(5)*", !"int addrspace(5)*"}
1434 !60 = !{i32 1, i32 1, i32 1}
1435 !61 = !{!"read_only", !"write_only", !"read_write"}
1436 !62 = !{!"image1d_t", !"image2d_t", !"image3d_t"}
1437 !70 = !{!"volatile", !"const restrict", !"pipe"}
1438 !80 = !{!"int addrspace(5)* addrspace(5)*"}
1439 !81 = !{i32 1}
1440 !82 = !{!"struct B"}
1441 !83 = !{!"global int addrspace(5)* __attribute__((ext_vector_type(2)))"}
1442 !84 = !{!"clk_event_t"}
1443 !opencl.ocl.version = !{!90}
1444 !90 = !{i32 2, i32 0}
1445 !91 = !{i32 0, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3}
1446 !92 = !{!"none", !"none", !"none", !"none", !"none", !"none", !"none"}
1447 !93 = !{!"long addrspace(5)*", !"char addrspace(5)*", !"char2 addrspace(5)*", !"char3 addrspace(5)*", !"char4 addrspace(5)*", !"char8 addrspace(5)*", !"char16 addrspace(5)*"}
1448 !94 = !{!"", !"", !"", !"", !"", !"", !""}
1449 !100 = !{!"1:1:4:%d\5Cn"}
1450 !101 = !{!"2:1:8:%g\5Cn"}
1451 !110 = !{!"__block_literal"}
1453 ; PARSER: AMDGPU HSA Metadata Parser Test: PASS