1// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK %s
2// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK %s
3// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK %s
4
5// CHECK:      	.amdgpu_metadata
6// CHECK:      amdhsa.kernels:
7// CHECK:        - .group_segment_fixed_size: 24
8// CHECK:          .kernarg_segment_align: 16
9// CHECK:          .kernarg_segment_size: 24
10// CHECK:          .max_flat_workgroup_size: 256
11// CHECK:          .name:           test_kernel
12// CHECK:          .private_segment_fixed_size: 16
13// CHECK:          .sgpr_count:     40
14// CHECK:          .sgpr_spill_count: 1
15// CHECK:          .symbol:         'test_kernel@kd'
16// CHECK:          .uses_dynamic_stack: true
17// CHECK:          .vgpr_count:     14
18// CHECK:          .vgpr_spill_count: 1
19// CHECK:          .wavefront_size: 64
20// CHECK:      amdhsa.version:
21// CHECK-NEXT:   - 1
22// CHECK-NEXT:   - 0
23.amdgpu_metadata
24  amdhsa.version:
25    - 1
26    - 0
27  amdhsa.printf:
28    - '1:1:4:%d\n'
29    - '2:1:8:%g\n'
30  amdhsa.kernels:
31    - .name:            test_kernel
32      .symbol:      test_kernel@kd
33      .kernarg_segment_size:      24
34      .group_segment_fixed_size:   24
35      .private_segment_fixed_size: 16
36      .uses_dynamic_stack: true
37      .kernarg_segment_align:     16
38      .wavefront_size:           64
39      .max_flat_workgroup_size:    256
40      .sgpr_count:               40
41      .vgpr_count:               14
42      .sgpr_spill_count:         1
43      .vgpr_spill_count:         1
44.end_amdgpu_metadata
45