1// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s 2// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s 3// RUN: llvm-mc -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: - .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: 12// CHECK-NEXT: - 2 13// CHECK-NEXT: - 0 14// CHECK: .max_flat_workgroup_size: 256 15// CHECK: .name: test_kernel 16// CHECK: .private_segment_fixed_size: 32 17// CHECK: .reqd_workgroup_size: 18// CHECK-NEXT: - 1 19// CHECK-NEXT: - 2 20// CHECK-NEXT: - 4 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: 27// CHECK-NEXT: - 8 28// CHECK-NEXT: - 16 29// CHECK-NEXT: - 32 30// CHECK: amdhsa.printf: 31// CHECK: - '1:1:4:%d\n' 32// CHECK: - '2:1:8:%g\n' 33// CHECK: amdhsa.version: 34// CHECK-NEXT: - 1 35// CHECK-NEXT: - 0 36// CHECK: .end_amdgpu_metadata 37.amdgpu_metadata 38 amdhsa.version: 39 - 1 40 - 0 41 amdhsa.printf: 42 - '1:1:4:%d\n' 43 - '2:1:8:%g\n' 44 amdhsa.kernels: 45 - .name: test_kernel 46 .symbol: test_kernel@kd 47 .language: OpenCL C 48 .language_version: 49 - 2 50 - 0 51 .kernarg_segment_size: 8 52 .group_segment_fixed_size: 16 53 .private_segment_fixed_size: 32 54 .kernarg_segment_align: 64 55 .wavefront_size: 128 56 .sgpr_count: 14 57 .vgpr_count: 40 58 .max_flat_workgroup_size: 256 59 .reqd_workgroup_size: 60 - 1 61 - 2 62 - 4 63 .workgroup_size_hint: 64 - 8 65 - 16 66 - 32 67 .vec_type_hint: int 68.end_amdgpu_metadata 69