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