97 lines · plain
1// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK %s2// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK %s3// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK %s4 5; CHECK: .amdgpu_metadata6; CHECK: amdhsa.kernels: 7; CHECK-NEXT: - .args: 8; CHECK-NEXT: - .offset: 19; CHECK-NEXT: .size: 110; CHECK-NEXT: .type_name: char11; CHECK-NEXT: .value_kind: by_value12; CHECK-NEXT: .value_type: i813; CHECK-NEXT: - .offset: 814; CHECK-NEXT: .size: 815; CHECK-NEXT: .value_kind: hidden_global_offset_x16; CHECK-NEXT: .value_type: i6417; CHECK-NEXT: - .offset: 818; CHECK-NEXT: .size: 819; CHECK-NEXT: .value_kind: hidden_global_offset_y20; CHECK-NEXT: .value_type: i6421; CHECK-NEXT: - .offset: 822; CHECK-NEXT: .size: 823; CHECK-NEXT: .value_kind: hidden_global_offset_z24; CHECK-NEXT: .value_type: i6425; CHECK-NEXT: - .address_space: global26; CHECK-NEXT: .offset: 827; CHECK-NEXT: .size: 828; CHECK-NEXT: .value_kind: hidden_printf_buffer29; CHECK-NEXT: .value_type: i830; CHECK-NEXT: .group_segment_fixed_size: 1631; CHECK-NEXT: .kernarg_segment_align: 6432; CHECK-NEXT: .kernarg_segment_size: 833; CHECK-NEXT: .language: OpenCL C34; CHECK-NEXT: .language_version: 35; CHECK-NEXT: - 236; CHECK-NEXT: - 037; CHECK-NEXT: .max_flat_workgroup_size: 25638; CHECK-NEXT: .name: test_kernel39; CHECK-NEXT: .private_segment_fixed_size: 3240; CHECK-NEXT: .sgpr_count: 1441; CHECK-NEXT: .symbol: 'test_kernel@kd'42; CHECK-NEXT: .vgpr_count: 4043; CHECK-NEXT: .wavefront_size: 12844; CHECK-NEXT: amdhsa.printf: 45; CHECK-NEXT: - '1:1:4:%d\n'46; CHECK-NEXT: - '2:1:8:%g\n'47; CHECK-NEXT: amdhsa.version: 48; CHECK-NEXT: - 149; CHECK-NEXT: - 050; CHECK: .end_amdgpu_metadata51.amdgpu_metadata52 amdhsa.version:53 - 154 - 055 amdhsa.printf:56 - '1:1:4:%d\n'57 - '2:1:8:%g\n'58 amdhsa.kernels:59 - .name: test_kernel60 .symbol: test_kernel@kd61 .language: OpenCL C62 .language_version:63 - 264 - 065 .kernarg_segment_size: 866 .group_segment_fixed_size: 1667 .private_segment_fixed_size: 3268 .kernarg_segment_align: 6469 .wavefront_size: 12870 .sgpr_count: 1471 .vgpr_count: 4072 .max_flat_workgroup_size: 25673 .args:74 - .type_name: char75 .size: 176 .offset: 177 .value_kind: by_value78 .value_type: i879 - .size: 880 .offset: 881 .value_kind: hidden_global_offset_x82 .value_type: i6483 - .size: 884 .offset: 885 .value_kind: hidden_global_offset_y86 .value_type: i6487 - .size: 888 .offset: 889 .value_kind: hidden_global_offset_z90 .value_type: i6491 - .size: 892 .offset: 893 .value_kind: hidden_printf_buffer94 .value_type: i895 .address_space: global96.end_amdgpu_metadata97