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
5 // CHECK
: .amdgpu_metadata
6 // CHECK
: amdhsa.kernels
:
7 // CHECK
: - .group_segment_fixed_size: 16
8 // CHECK
: .kernarg_segment_align: 64
9 // CHECK
: .kernarg_segment_size: 8
10 // CHECK
: .language: OpenCL C
11 // CHECK
: .language_version:
14 // CHECK
: .max_flat_workgroup_size: 256
15 // CHECK
: .name: test_kernel
16 // CHECK
: .private_segment_fixed_size: 32
17 // CHECK
: .reqd_workgroup_size:
21 // CHECK
: .sgpr_count: 14
22 // CHECK
: .symbol: 'test_kernel@kd'
23 // CHECK
: .vec_type_hint: int
24 // CHECK
: .vgpr_count: 40
25 // CHECK
: .wavefront_size: 128
26 // CHECK
: .workgroup_size_hint:
30 // CHECK
: amdhsa.printf
:
31 // CHECK
: - '1:1:4:%d\n'
32 // CHECK
: - '2:1:8:%g\n'
33 // CHECK
: amdhsa.version
:
36 // CHECK
: .end_amdgpu_metadata
46 .symbol: test_kernel@kd
51 .kernarg_segment_size: 8
52 .group_segment_fixed_size: 16
53 .private_segment_fixed_size: 32
54 .kernarg_segment_align: 64
58 .max_flat_workgroup_size: 256