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