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:      - .max_flat_workgroup_size: 256
8// CHECK:        .wavefront_size: 128
9// CHECK:        .symbol:      'test_kernel@kd'
10// CHECK:        .reqd_workgroup_size:
11// CHECK-NEXT:     - 1
12// CHECK-NEXT:     - 2
13// CHECK-NEXT:     - 4
14// CHECK:        .kernarg_segment_size: 8
15// CHECK:        .private_segment_fixed_size: 32
16// CHECK:        .workgroup_size_hint:
17// CHECK-NEXT:     - 8
18// CHECK-NEXT:     - 16
19// CHECK-NEXT:     - 32
20// CHECK:        .name:            test_kernel
21// CHECK:        .language:        OpenCL C
22// CHECK:        .sgpr_count:      14
23// CHECK:        .kernarg_segment_align: 64
24// CHECK:        .vgpr_count:      40
25// CHECK:        .language_version:
26// CHECK-NEXT:     - 2
27// CHECK-NEXT:     - 0
28// CHECK:        .vec_type_hint:       int
29// CHECK:    amdhsa.version:
30// CHECK-NEXT: - 1
31// CHECK-NEXT: - 0
32// CHECK:    amdhsa.printf:
33// CHECK:      - '1:1:4:%d\n'
34// CHECK:      - '2:1:8:%g\n'
35// CHECK: .end_amdgpu_metadata
36.amdgpu_metadata
37  amdhsa.version:
38    - 1
39    - 0
40  amdhsa.printf:
41    - '1:1:4:%d\n'
42    - '2:1:8:%g\n'
43  amdhsa.kernels:
44    - .name:            test_kernel
45      .symbol:      test_kernel@kd
46      .language:        OpenCL C
47      .language_version:
48        - 2
49        - 0
50      .kernarg_segment_size: 8
51      .group_segment_fixed_size: 16
52      .private_segment_fixed_size: 32
53      .kernarg_segment_align: 64
54      .wavefront_size: 128
55      .sgpr_count: 14
56      .vgpr_count: 40
57      .max_flat_workgroup_size: 256
58      .reqd_workgroup_size:
59        - 1
60        - 2
61        - 4
62      .workgroup_size_hint:
63        - 8
64        - 16
65        - 32
66      .vec_type_hint:       int
67.end_amdgpu_metadata
68