Gitiles
Code Review Sign In
gerrit-public.fairphone.software / toolchain / llvm-project / e29e30b1397f3e50f3487491f8a77ae08e4e3471 / . / llvm / test / CodeGen / AMDGPU / hsa-metadata-from-llvm-ir-full-v3.ll
blob: edcd45a3fbfa066ba8b55bb7a302b27ac00cea8f [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
47; CHECK: .language: OpenCL C
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]48; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]49; CHECK-NEXT: - 2
50; CHECK-NEXT: - 0
51; CHECK: .name: test_char
52; CHECK: .symbol: test_char.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]53define amdgpu_kernel void @test_char(i8 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]59; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]89; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]90; CHECK-NEXT: - 2
91; CHECK-NEXT: - 0
92; CHECK: .name: test_ushort2
93; CHECK: .symbol: test_ushort2.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]94define amdgpu_kernel void @test_ushort2(<2 x i16> %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]100; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]125; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]126; CHECK-NEXT: - 2
127; CHECK-NEXT: - 0
128; CHECK: .name: test_int3
129; CHECK: .symbol: test_int3.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]130define amdgpu_kernel void @test_int3(<3 x i32> %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]136; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]161; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]162; CHECK-NEXT: - 2
163; CHECK-NEXT: - 0
164; CHECK: .name: test_ulong4
165; CHECK: .symbol: test_ulong4.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]166define amdgpu_kernel void @test_ulong4(<4 x i64> %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]172; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]197; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]198; CHECK-NEXT: - 2
199; CHECK-NEXT: - 0
200; CHECK: .name: test_half8
201; CHECK: .symbol: test_half8.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]202define amdgpu_kernel void @test_half8(<8 x half> %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]208; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]233; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]234; CHECK-NEXT: - 2
235; CHECK-NEXT: - 0
236; CHECK: .name: test_float16
237; CHECK: .symbol: test_float16.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]238define amdgpu_kernel void @test_float16(<16 x float> %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]244; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]269; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]270; CHECK-NEXT: - 2
271; CHECK-NEXT: - 0
272; CHECK: .name: test_double16
273; CHECK: .symbol: test_double16.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]274define amdgpu_kernel void @test_double16(<16 x double> %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]280; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]306; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]307; CHECK-NEXT: - 2
308; CHECK-NEXT: - 0
309; CHECK: .name: test_pointer
310; CHECK: .symbol: test_pointer.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]311define amdgpu_kernel void @test_pointer(i32 addrspace(1)* %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]317; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]343; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]344; CHECK-NEXT: - 2
345; CHECK-NEXT: - 0
346; CHECK: .name: test_image
347; CHECK: .symbol: test_image.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]348define amdgpu_kernel void @test_image(%opencl.image2d_t addrspace(1)* %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]354; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]379; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]380; CHECK-NEXT: - 2
381; CHECK-NEXT: - 0
382; CHECK: .name: test_sampler
383; CHECK: .symbol: test_sampler.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]384define amdgpu_kernel void @test_sampler(i32 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]390; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]416; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]417; CHECK-NEXT: - 2
418; CHECK-NEXT: - 0
419; CHECK: .name: test_queue
420; CHECK: .symbol: test_queue.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]421define amdgpu_kernel void @test_queue(%opencl.queue_t addrspace(1)* %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]427; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]453; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]454; CHECK-NEXT: - 2
455; CHECK-NEXT: - 0
456; CHECK: .name: test_struct
457; CHECK: .symbol: test_struct.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]458define amdgpu_kernel void @test_struct(%struct.A addrspace(5)* byval %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]464; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]489; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]490; CHECK-NEXT: - 2
491; CHECK-NEXT: - 0
492; CHECK: .name: test_i128
493; CHECK: .symbol: test_i128.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]494define amdgpu_kernel void @test_i128(i128 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]500; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]537; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]538; CHECK-NEXT: - 2
539; CHECK-NEXT: - 0
540; CHECK: .name: test_multi_arg
541; CHECK: .symbol: test_multi_arg.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]542define 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]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]548; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]589; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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,
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]596 i32 addrspace(3)* %l) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]602; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]646; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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,
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]653 %opencl.pipe_t addrspace(1)* %c) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]659; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]702; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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,
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]709 %opencl.image3d_t addrspace(1)* %rw) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]715; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]740; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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 Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]746define amdgpu_kernel void @test_vec_type_hint_half(i32 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]752; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]777; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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 Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]783define amdgpu_kernel void @test_vec_type_hint_float(i32 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]789; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]814; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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 Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]820define amdgpu_kernel void @test_vec_type_hint_double(i32 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]826; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]851; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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 Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]857define amdgpu_kernel void @test_vec_type_hint_char(i32 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]863; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]888; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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 Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]894define amdgpu_kernel void @test_vec_type_hint_short(i32 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]900; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]925; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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 Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]931define amdgpu_kernel void @test_vec_type_hint_long(i32 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]937; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]962; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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 Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]968define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]974; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]999; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1000; CHECK-NEXT: - 2
1001; CHECK-NEXT: - 0
1002; CHECK: .name: test_reqd_wgs_vec_type_hint
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1003; CHECK: .reqd_workgroup_size:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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 Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1009define amdgpu_kernel void @test_reqd_wgs_vec_type_hint(i32 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1016; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1041; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1047; CHECK: .workgroup_size_hint:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1048; CHECK-NEXT: - 8
1049; CHECK-NEXT: - 16
1050; CHECK-NEXT: - 32
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1051define amdgpu_kernel void @test_wgs_hint_vec_type_hint(i32 %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1058; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1084; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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 Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1089define 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]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1095; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1121; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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 Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1126define 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]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1132; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1157; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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 Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1162define 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]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1168; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1194; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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(
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1200 %opencl.clk_event_t addrspace(1)* %a) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1206; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]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
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1262; CHECK-NEXT: - .address_space: local
1263; CHECK-NEXT: .name: h
1264; CHECK-NEXT: .offset: 32
1265; CHECK-NEXT: .pointee_align: 1
1266; CHECK-NEXT: .size: 4
1267; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1268; CHECK-NEXT: .value_type: struct
1269; CHECK-NEXT: - .offset: 40
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1270; CHECK-NEXT: .size: 8
1271; CHECK-NEXT: .value_kind: hidden_global_offset_x
1272; CHECK-NEXT: .value_type: i64
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1273; CHECK-NEXT: - .offset: 48
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1274; CHECK-NEXT: .size: 8
1275; CHECK-NEXT: .value_kind: hidden_global_offset_y
1276; CHECK-NEXT: .value_type: i64
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1277; CHECK-NEXT: - .offset: 56
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1278; CHECK-NEXT: .size: 8
1279; CHECK-NEXT: .value_kind: hidden_global_offset_z
1280; CHECK-NEXT: .value_type: i64
1281; CHECK-NEXT: - .address_space: global
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1282; CHECK-NEXT: .offset: 64
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1283; CHECK-NEXT: .size: 8
1284; CHECK-NEXT: .value_kind: hidden_printf_buffer
1285; CHECK-NEXT: .value_type: i8
1286; CHECK: .language: OpenCL C
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1287; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1288; CHECK-NEXT: - 2
1289; CHECK-NEXT: - 0
1290; CHECK: .name: test_pointee_align
1291; CHECK: .symbol: test_pointee_align.kd
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]1292define amdgpu_kernel void @test_pointee_align(i64 addrspace(1)* %a,
1293 i8 addrspace(3)* %b,
1294 <2 x i8> addrspace(3)* %c,
1295 <3 x i8> addrspace(3)* %d,
1296 <4 x i8> addrspace(3)* %e,
1297 <8 x i8> addrspace(3)* %f,
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1298 <16 x i8> addrspace(3)* %g,
1299 {} addrspace(3)* %h) #0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]1300 !kernel_arg_addr_space !91 !kernel_arg_access_qual !92 !kernel_arg_type !93
1301 !kernel_arg_base_type !93 !kernel_arg_type_qual !94 {
1302 ret void
1303}
1304
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1305; CHECK: - .args:
1306; CHECK-NEXT: - .address_space: global
1307; CHECK-NEXT: .name: a
1308; CHECK-NEXT: .offset: 0
1309; CHECK-NEXT: .size: 8
1310; CHECK-NEXT: .type_name: 'long addrspace(5)*'
1311; CHECK-NEXT: .value_kind: global_buffer
1312; CHECK-NEXT: .value_type: i64
1313; CHECK-NEXT: - .address_space: local
1314; CHECK-NEXT: .name: b
1315; CHECK-NEXT: .offset: 8
1316; CHECK-NEXT: .pointee_align: 8
1317; CHECK-NEXT: .size: 4
1318; CHECK-NEXT: .type_name: 'char addrspace(5)*'
1319; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1320; CHECK-NEXT: .value_type: i8
1321; CHECK-NEXT: - .address_space: local
1322; CHECK-NEXT: .name: c
1323; CHECK-NEXT: .offset: 12
1324; CHECK-NEXT: .pointee_align: 32
1325; CHECK-NEXT: .size: 4
1326; CHECK-NEXT: .type_name: 'char2 addrspace(5)*'
1327; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1328; CHECK-NEXT: .value_type: i8
1329; CHECK-NEXT: - .address_space: local
1330; CHECK-NEXT: .name: d
1331; CHECK-NEXT: .offset: 16
1332; CHECK-NEXT: .pointee_align: 64
1333; CHECK-NEXT: .size: 4
1334; CHECK-NEXT: .type_name: 'char3 addrspace(5)*'
1335; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1336; CHECK-NEXT: .value_type: i8
1337; CHECK-NEXT: - .address_space: local
1338; CHECK-NEXT: .name: e
1339; CHECK-NEXT: .offset: 20
1340; CHECK-NEXT: .pointee_align: 256
1341; CHECK-NEXT: .size: 4
1342; CHECK-NEXT: .type_name: 'char4 addrspace(5)*'
1343; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1344; CHECK-NEXT: .value_type: i8
1345; CHECK-NEXT: - .address_space: local
1346; CHECK-NEXT: .name: f
1347; CHECK-NEXT: .offset: 24
1348; CHECK-NEXT: .pointee_align: 128
1349; CHECK-NEXT: .size: 4
1350; CHECK-NEXT: .type_name: 'char8 addrspace(5)*'
1351; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1352; CHECK-NEXT: .value_type: i8
1353; CHECK-NEXT: - .address_space: local
1354; CHECK-NEXT: .name: g
1355; CHECK-NEXT: .offset: 28
1356; CHECK-NEXT: .pointee_align: 1024
1357; CHECK-NEXT: .size: 4
1358; CHECK-NEXT: .type_name: 'char16 addrspace(5)*'
1359; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1360; CHECK-NEXT: .value_type: i8
1361; CHECK-NEXT: - .address_space: local
1362; CHECK-NEXT: .name: h
1363; CHECK-NEXT: .offset: 32
1364; CHECK-NEXT: .pointee_align: 16
1365; CHECK-NEXT: .size: 4
1366; CHECK-NEXT: .value_kind: dynamic_shared_pointer
1367; CHECK-NEXT: .value_type: struct
1368; CHECK-NEXT: - .offset: 40
1369; CHECK-NEXT: .size: 8
1370; CHECK-NEXT: .value_kind: hidden_global_offset_x
1371; CHECK-NEXT: .value_type: i64
1372; CHECK-NEXT: - .offset: 48
1373; CHECK-NEXT: .size: 8
1374; CHECK-NEXT: .value_kind: hidden_global_offset_y
1375; CHECK-NEXT: .value_type: i64
1376; CHECK-NEXT: - .offset: 56
1377; CHECK-NEXT: .size: 8
1378; CHECK-NEXT: .value_kind: hidden_global_offset_z
1379; CHECK-NEXT: .value_type: i64
1380; CHECK-NEXT: - .address_space: global
1381; CHECK-NEXT: .offset: 64
1382; CHECK-NEXT: .size: 8
1383; CHECK-NEXT: .value_kind: hidden_printf_buffer
1384; CHECK-NEXT: .value_type: i8
1385; CHECK: .language: OpenCL C
1386; CHECK-NEXT: .language_version:
1387; CHECK-NEXT: - 2
1388; CHECK-NEXT: - 0
1389; CHECK: .name: test_pointee_align_attribute
1390; CHECK: .symbol: test_pointee_align_attribute.kd
1391define amdgpu_kernel void @test_pointee_align_attribute(i64 addrspace(1)* align 16 %a,
1392 i8 addrspace(3)* align 8 %b,
1393 <2 x i8> addrspace(3)* align 32 %c,
1394 <3 x i8> addrspace(3)* align 64 %d,
1395 <4 x i8> addrspace(3)* align 256 %e,
1396 <8 x i8> addrspace(3)* align 128 %f,
1397 <16 x i8> addrspace(3)* align 1024 %g,
1398 {} addrspace(3)* align 16 %h) #0
1399 !kernel_arg_addr_space !91 !kernel_arg_access_qual !92 !kernel_arg_type !93
1400 !kernel_arg_base_type !93 !kernel_arg_type_qual !94 {
1401 ret void
1402}
1403; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1404; CHECK-NEXT: - .name: arg
1405; CHECK-NEXT: .offset: 0
1406; CHECK-NEXT: .size: 25
1407; CHECK-NEXT: .type_name: __block_literal
1408; CHECK-NEXT: .value_kind: by_value
1409; CHECK-NEXT: .value_type: struct
1410; CHECK-NEXT: - .offset: 32
1411; CHECK-NEXT: .size: 8
1412; CHECK-NEXT: .value_kind: hidden_global_offset_x
1413; CHECK-NEXT: .value_type: i64
1414; CHECK-NEXT: - .offset: 40
1415; CHECK-NEXT: .size: 8
1416; CHECK-NEXT: .value_kind: hidden_global_offset_y
1417; CHECK-NEXT: .value_type: i64
1418; CHECK-NEXT: - .offset: 48
1419; CHECK-NEXT: .size: 8
1420; CHECK-NEXT: .value_kind: hidden_global_offset_z
1421; CHECK-NEXT: .value_type: i64
1422; CHECK-NEXT: - .address_space: global
1423; CHECK-NEXT: .offset: 56
1424; CHECK-NEXT: .size: 8
1425; CHECK-NEXT: .value_kind: hidden_printf_buffer
1426; CHECK-NEXT: .value_type: i8
1427; CHECK: .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle
1428; CHECK: .language: OpenCL C
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1429; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1430; CHECK-NEXT: - 2
1431; CHECK-NEXT: - 0
1432; CHECK: .name: __test_block_invoke_kernel
1433; CHECK: .symbol: __test_block_invoke_kernel.kd
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]1434define amdgpu_kernel void @__test_block_invoke_kernel(
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1435 <{ i32, i32, i8*, i8 addrspace(1)*, i8 }> %arg) #1
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]1436 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !110
1437 !kernel_arg_base_type !110 !kernel_arg_type_qual !4 {
1438 ret void
1439}
1440
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1441; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1442; CHECK-NEXT: - .name: a
1443; CHECK-NEXT: .offset: 0
1444; CHECK-NEXT: .size: 1
1445; CHECK-NEXT: .type_name: char
1446; CHECK-NEXT: .value_kind: by_value
1447; CHECK-NEXT: .value_type: i8
1448; CHECK-NEXT: - .offset: 8
1449; CHECK-NEXT: .size: 8
1450; CHECK-NEXT: .value_kind: hidden_global_offset_x
1451; CHECK-NEXT: .value_type: i64
1452; CHECK-NEXT: - .offset: 16
1453; CHECK-NEXT: .size: 8
1454; CHECK-NEXT: .value_kind: hidden_global_offset_y
1455; CHECK-NEXT: .value_type: i64
1456; CHECK-NEXT: - .offset: 24
1457; CHECK-NEXT: .size: 8
1458; CHECK-NEXT: .value_kind: hidden_global_offset_z
1459; CHECK-NEXT: .value_type: i64
1460; CHECK-NEXT: - .address_space: global
1461; CHECK-NEXT: .offset: 32
1462; CHECK-NEXT: .size: 8
1463; CHECK-NEXT: .value_kind: hidden_printf_buffer
1464; CHECK-NEXT: .value_type: i8
1465; CHECK-NEXT: - .address_space: global
1466; CHECK-NEXT: .offset: 40
1467; CHECK-NEXT: .size: 8
1468; CHECK-NEXT: .value_kind: hidden_default_queue
1469; CHECK-NEXT: .value_type: i8
1470; CHECK-NEXT: - .address_space: global
1471; CHECK-NEXT: .offset: 48
1472; CHECK-NEXT: .size: 8
1473; CHECK-NEXT: .value_kind: hidden_completion_action
1474; CHECK-NEXT: .value_type: i8
1475; CHECK: .language: OpenCL C
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1476; CHECK-NEXT: .language_version:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1477; CHECK-NEXT: - 2
1478; CHECK-NEXT: - 0
1479; CHECK: .name: test_enqueue_kernel_caller
1480; CHECK: .symbol: test_enqueue_kernel_caller.kd
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1481define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #2
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]1482 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
1483 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
1484 ret void
1485}
1486
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1487; CHECK: - .args:
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1488; CHECK-NEXT: - .name: ptr
1489; CHECK-NEXT: .offset: 0
1490; CHECK-NEXT: .size: 8
1491; CHECK-NEXT: .value_kind: global_buffer
1492; CHECK-NEXT: .value_type: i32
1493; CHECK: .name: unknown_addrspace_kernarg
1494; CHECK: .symbol: unknown_addrspace_kernarg.kd
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]1495define amdgpu_kernel void @unknown_addrspace_kernarg(i32 addrspace(12345)* %ptr) #0 {
1496 ret void
1497}
1498
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]1499; CHECK: amdhsa.printf:
1500; CHECK-NEXT: - '1:1:4:%d\n'
1501; CHECK-NEXT: - '2:1:8:%g\n'
Tim Renoufed0b9af2019-03-13 18:55:50 +0000[diff] [blame]1502; CHECK: amdhsa.version:
1503; CHECK-NEXT: - 1
1504; CHECK-NEXT: - 0
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]1505
Scott Linder3eed9612019-04-23 14:31:17 +0000[diff] [blame]1506attributes #0 = { "amdgpu-implicitarg-num-bytes"="48" }
1507attributes #1 = { "amdgpu-implicitarg-num-bytes"="48" "runtime-handle"="__test_block_invoke_kernel_runtime_handle" }
1508attributes #2 = { "amdgpu-implicitarg-num-bytes"="48" "calls-enqueue-kernel" }
Scott Linderf5b36e52018-12-12 19:39:27 +0000[diff] [blame]1509
1510!llvm.printf.fmts = !{!100, !101}
1511
1512!1 = !{i32 0}
1513!2 = !{!"none"}
1514!3 = !{!"int"}
1515!4 = !{!""}
1516!5 = !{i32 undef, i32 1}
1517!6 = !{i32 1, i32 2, i32 4}
1518!7 = !{<4 x i32> undef, i32 0}
1519!8 = !{i32 8, i32 16, i32 32}
1520!9 = !{!"char"}
1521!10 = !{!"ushort2"}
1522!11 = !{!"int3"}
1523!12 = !{!"ulong4"}
1524!13 = !{!"half8"}
1525!14 = !{!"float16"}
1526!15 = !{!"double16"}
1527!16 = !{!"int addrspace(5)*"}
1528!17 = !{!"image2d_t"}
1529!18 = !{!"sampler_t"}
1530!19 = !{!"queue_t"}
1531!20 = !{!"struct A"}
1532!21 = !{!"i128"}
1533!22 = !{i32 0, i32 0, i32 0}
1534!23 = !{!"none", !"none", !"none"}
1535!24 = !{!"int", !"short2", !"char3"}
1536!25 = !{!"", !"", !""}
1537!26 = !{half undef, i32 1}
1538!27 = !{float undef, i32 1}
1539!28 = !{double undef, i32 1}
1540!29 = !{i8 undef, i32 1}
1541!30 = !{i16 undef, i32 1}
1542!31 = !{i64 undef, i32 1}
1543!32 = !{i32 addrspace(5)*undef, i32 1}
1544!50 = !{i32 1, i32 2, i32 3}
1545!51 = !{!"int addrspace(5)*", !"int addrspace(5)*", !"int addrspace(5)*"}
1546!60 = !{i32 1, i32 1, i32 1}
1547!61 = !{!"read_only", !"write_only", !"read_write"}
1548!62 = !{!"image1d_t", !"image2d_t", !"image3d_t"}
1549!70 = !{!"volatile", !"const restrict", !"pipe"}
1550!80 = !{!"int addrspace(5)* addrspace(5)*"}
1551!81 = !{i32 1}
1552!82 = !{!"struct B"}
1553!83 = !{!"global int addrspace(5)* __attribute__((ext_vector_type(2)))"}
1554!84 = !{!"clk_event_t"}
1555!opencl.ocl.version = !{!90}
1556!90 = !{i32 2, i32 0}
1557!91 = !{i32 0, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3}
1558!92 = !{!"none", !"none", !"none", !"none", !"none", !"none", !"none"}
1559!93 = !{!"long addrspace(5)*", !"char addrspace(5)*", !"char2 addrspace(5)*", !"char3 addrspace(5)*", !"char4 addrspace(5)*", !"char8 addrspace(5)*", !"char16 addrspace(5)*"}
1560!94 = !{!"", !"", !"", !"", !"", !"", !""}
1561!100 = !{!"1:1:4:%d\5Cn"}
1562!101 = !{!"2:1:8:%g\5Cn"}
1563!110 = !{!"__block_literal"}
1564
1565; PARSER: AMDGPU HSA Metadata Parser Test: PASS