1// RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s
2// RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s
3// RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s
4
5; CHECK:      	.amdgpu_metadata
6; CHECK:      amdhsa.kernels:
7; CHECK-NEXT:   - .args:
8; CHECK-NEXT:       - .offset:         1
9; CHECK-NEXT:         .size:           1
10; CHECK-NEXT:         .type_name:      char
11; CHECK-NEXT:         .value_kind:     by_value
12; CHECK-NEXT:         .value_type:     i8
13; CHECK-NEXT:       - .offset:         8
14; CHECK-NEXT:         .size:           8
15; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
16; CHECK-NEXT:         .value_type:     i64
17; CHECK-NEXT:       - .offset:         8
18; CHECK-NEXT:         .size:           8
19; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
20; CHECK-NEXT:         .value_type:     i64
21; CHECK-NEXT:       - .offset:         8
22; CHECK-NEXT:         .size:           8
23; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
24; CHECK-NEXT:         .value_type:     i64
25; CHECK-NEXT:       - .address_space:  global
26; CHECK-NEXT:         .offset:         8
27; CHECK-NEXT:         .size:           8
28; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
29; CHECK-NEXT:         .value_type:     i8
30; CHECK-NEXT:     .group_segment_fixed_size: 16
31; CHECK-NEXT:     .kernarg_segment_align: 64
32; CHECK-NEXT:     .kernarg_segment_size: 8
33; CHECK-NEXT:     .language:       OpenCL C
34; CHECK-NEXT:     .language_version:
35; CHECK-NEXT:       - 2
36; CHECK-NEXT:       - 0
37; CHECK-NEXT:     .max_flat_workgroup_size: 256
38; CHECK-NEXT:     .name:           test_kernel
39; CHECK-NEXT:     .private_segment_fixed_size: 32
40; CHECK-NEXT:     .sgpr_count:     14
41; CHECK-NEXT:     .symbol:         'test_kernel@kd'
42; CHECK-NEXT:     .vgpr_count:     40
43; CHECK-NEXT:     .wavefront_size: 128
44; CHECK-NEXT: amdhsa.printf:
45; CHECK-NEXT:   - '1:1:4:%d\n'
46; CHECK-NEXT:   - '2:1:8:%g\n'
47; CHECK-NEXT: amdhsa.version:
48; CHECK-NEXT:   - 1
49; CHECK-NEXT:   - 0
50; CHECK:      	.end_amdgpu_metadata
51.amdgpu_metadata
52  amdhsa.version:
53    - 1
54    - 0
55  amdhsa.printf:
56    - '1:1:4:%d\n'
57    - '2:1:8:%g\n'
58  amdhsa.kernels:
59    - .name:            test_kernel
60      .symbol:      test_kernel@kd
61      .language:        OpenCL C
62      .language_version:
63        - 2
64        - 0
65      .kernarg_segment_size: 8
66      .group_segment_fixed_size: 16
67      .private_segment_fixed_size: 32
68      .kernarg_segment_align: 64
69      .wavefront_size: 128
70      .sgpr_count: 14
71      .vgpr_count: 40
72      .max_flat_workgroup_size: 256
73      .args:
74        - .type_name:      char
75          .size:          1
76          .offset:         1
77          .value_kind:     by_value
78          .value_type:     i8
79        - .size:          8
80          .offset:         8
81          .value_kind:     hidden_global_offset_x
82          .value_type:     i64
83        - .size:          8
84          .offset:         8
85          .value_kind:     hidden_global_offset_y
86          .value_type:     i64
87        - .size:          8
88          .offset:         8
89          .value_kind:     hidden_global_offset_z
90          .value_type:     i64
91        - .size:          8
92          .offset:         8
93          .value_kind:     hidden_printf_buffer
94          .value_type:     i8
95          .address_space: global
96.end_amdgpu_metadata
97