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