297 lines · plain
1// RUN: llvm-mc -triple amdgcn-amd-amdhsa -mcpu=gfx1200 < %s | FileCheck --check-prefix=ASM %s2// RUN: llvm-mc -triple amdgcn-amd-amdhsa -mcpu=gfx1200 -filetype=obj < %s > %t3// RUN: llvm-readelf -S -r -s %t | FileCheck --check-prefix=READOBJ %s4// RUN: llvm-objdump -s -j .rodata %t | FileCheck --check-prefix=OBJDUMP %s5 6// READOBJ: Section Headers7// READOBJ: .text PROGBITS {{[0-9a-f]+}} {{[0-9a-f]+}} {{[0-9a-f]+}} {{[0-9]+}} AX {{[0-9]+}} {{[0-9]+}} 2568// READOBJ: .rodata PROGBITS {{[0-9a-f]+}} {{[0-9a-f]+}} 000100 {{[0-9]+}} A {{[0-9]+}} {{[0-9]+}} 649 10// READOBJ: Relocation section '.rela.rodata' at offset11// READOBJ: 0000000000000010 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 1012// READOBJ: 0000000000000050 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 11013// READOBJ: 0000000000000090 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 21014// READOBJ: 00000000000000d0 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 31015 16// READOBJ: Symbol table '.symtab' contains {{[0-9]+}} entries:17// READOBJ: 0000000000000000 0 FUNC LOCAL PROTECTED 2 minimal18// READOBJ-NEXT: 0000000000000100 0 FUNC LOCAL PROTECTED 2 complete19// READOBJ-NEXT: 0000000000000200 0 FUNC LOCAL PROTECTED 2 special_sgpr20// READOBJ-NEXT: 0000000000000300 0 FUNC LOCAL PROTECTED 2 disabled_user_sgpr21// READOBJ-NEXT: 0000000000000000 64 OBJECT LOCAL DEFAULT 3 minimal.kd22// READOBJ-NEXT: 0000000000000040 64 OBJECT LOCAL DEFAULT 3 complete.kd23// READOBJ-NEXT: 0000000000000080 64 OBJECT LOCAL DEFAULT 3 special_sgpr.kd24// READOBJ-NEXT: 00000000000000c0 64 OBJECT LOCAL DEFAULT 3 disabled_user_sgpr.kd25 26// OBJDUMP: Contents of section .rodata27// Note, relocation for KERNEL_CODE_ENTRY_BYTE_OFFSET is not resolved here.28// minimal29// OBJDUMP-NEXT: 0000 00000000 00000000 00000000 0000000030// OBJDUMP-NEXT: 0010 00000000 00000000 00000000 0000000031// OBJDUMP-NEXT: 0020 00000000 00000000 00000000 0000000032// OBJDUMP-NEXT: 0030 00000ce0 80000000 00040000 0000000033// complete34// OBJDUMP-NEXT: 0040 01000000 01000000 08000000 0000000035// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 0000000036// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 f00f000037// OBJDUMP-NEXT: 0070 015021e4 1f0f007f 5e040000 0000000038// special_sgpr39// OBJDUMP-NEXT: 0080 00000000 00000000 00000000 0000000040// OBJDUMP-NEXT: 0090 00000000 00000000 00000000 0000000041// OBJDUMP-NEXT: 00a0 00000000 00000000 00000000 0000000042// OBJDUMP-NEXT: 00b0 000000e0 80000000 00040000 0000000043// disabled_user_sgpr44// OBJDUMP-NEXT: 00c0 00000000 00000000 00000000 0000000045// OBJDUMP-NEXT: 00d0 00000000 00000000 00000000 0000000046// OBJDUMP-NEXT: 00e0 00000000 00000000 00000000 0000000047// OBJDUMP-NEXT: 00f0 00000ce0 80000000 00040000 0000000048 49.text50 51.amdgcn_target "amdgcn-amd-amdhsa--gfx1200"52// ASM: .amdgcn_target "amdgcn-amd-amdhsa--gfx1200"53 54.amdhsa_code_object_version 455// ASM: .amdhsa_code_object_version 456 57.p2align 858.type minimal,@function59minimal:60 s_endpgm61 62.p2align 863.type complete,@function64complete:65 s_endpgm66 67.p2align 868.type special_sgpr,@function69special_sgpr:70 s_endpgm71 72.p2align 873.type disabled_user_sgpr,@function74disabled_user_sgpr:75 s_endpgm76 77.rodata78// ASM: .rodata79 80// Test that only specifying required directives is allowed, and that defaulted81// values are omitted.82.p2align 683.amdhsa_kernel minimal84 .amdhsa_next_free_vgpr 085 .amdhsa_next_free_sgpr 086.end_amdhsa_kernel87 88// ASM: .amdhsa_kernel minimal89// ASM: .amdhsa_next_free_vgpr 090// ASM-NEXT: .amdhsa_next_free_sgpr 091// ASM: .end_amdhsa_kernel92 93// Test that we can specify all available directives with non-default values.94.p2align 695.amdhsa_kernel complete96 .amdhsa_group_segment_fixed_size 197 .amdhsa_private_segment_fixed_size 198 .amdhsa_kernarg_size 899 .amdhsa_user_sgpr_count 15100 .amdhsa_user_sgpr_dispatch_ptr 1101 .amdhsa_user_sgpr_queue_ptr 1102 .amdhsa_user_sgpr_kernarg_segment_ptr 1103 .amdhsa_user_sgpr_dispatch_id 1104 .amdhsa_user_sgpr_private_segment_size 1105 .amdhsa_wavefront_size32 1106 .amdhsa_enable_private_segment 1107 .amdhsa_system_sgpr_workgroup_id_x 0108 .amdhsa_system_sgpr_workgroup_id_y 1109 .amdhsa_system_sgpr_workgroup_id_z 1110 .amdhsa_system_sgpr_workgroup_info 1111 .amdhsa_system_vgpr_workitem_id 1112 .amdhsa_next_free_vgpr 9113 .amdhsa_next_free_sgpr 27114 .amdhsa_reserve_vcc 0115 .amdhsa_float_round_mode_32 1116 .amdhsa_float_round_mode_16_64 1117 .amdhsa_float_denorm_mode_32 1118 .amdhsa_float_denorm_mode_16_64 0119 .amdhsa_fp16_overflow 1120 .amdhsa_workgroup_processor_mode 1121 .amdhsa_memory_ordered 1122 .amdhsa_forward_progress 1123 .amdhsa_inst_pref_size 255124 .amdhsa_round_robin_scheduling 1125 .amdhsa_exception_fp_ieee_invalid_op 1126 .amdhsa_exception_fp_denorm_src 1127 .amdhsa_exception_fp_ieee_div_zero 1128 .amdhsa_exception_fp_ieee_overflow 1129 .amdhsa_exception_fp_ieee_underflow 1130 .amdhsa_exception_fp_ieee_inexact 1131 .amdhsa_exception_int_div_zero 1132.end_amdhsa_kernel133 134// ASM: .amdhsa_kernel complete135// ASM-NEXT: .amdhsa_group_segment_fixed_size 1136// ASM-NEXT: .amdhsa_private_segment_fixed_size 1137// ASM-NEXT: .amdhsa_kernarg_size 8138// ASM-NEXT: .amdhsa_user_sgpr_count 15139// ASM-NEXT: .amdhsa_user_sgpr_dispatch_ptr 1140// ASM-NEXT: .amdhsa_user_sgpr_queue_ptr 1141// ASM-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 1142// ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1143// ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1144// ASM-NEXT: .amdhsa_wavefront_size32 1145// ASM-NEXT: .amdhsa_enable_private_segment 1146// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_x 0147// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_y 1148// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z 1149// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info 1150// ASM-NEXT: .amdhsa_system_vgpr_workitem_id 1151// ASM-NEXT: .amdhsa_next_free_vgpr 9152// ASM-NEXT: .amdhsa_next_free_sgpr 27153// ASM-NEXT: .amdhsa_reserve_vcc 0154// ASM-NEXT: .amdhsa_float_round_mode_32 1155// ASM-NEXT: .amdhsa_float_round_mode_16_64 1156// ASM-NEXT: .amdhsa_float_denorm_mode_32 1157// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 0158// ASM-NEXT: .amdhsa_fp16_overflow 1159// ASM-NEXT: .amdhsa_workgroup_processor_mode 1160// ASM-NEXT: .amdhsa_memory_ordered 1161// ASM-NEXT: .amdhsa_forward_progress 1162// ASM-NEXT: .amdhsa_inst_pref_size 255163// ASM-NEXT: .amdhsa_round_robin_scheduling 1164// ASM-NEXT: .amdhsa_exception_fp_ieee_invalid_op 1165// ASM-NEXT: .amdhsa_exception_fp_denorm_src 1166// ASM-NEXT: .amdhsa_exception_fp_ieee_div_zero 1167// ASM-NEXT: .amdhsa_exception_fp_ieee_overflow 1168// ASM-NEXT: .amdhsa_exception_fp_ieee_underflow 1169// ASM-NEXT: .amdhsa_exception_fp_ieee_inexact 1170// ASM-NEXT: .amdhsa_exception_int_div_zero 1171// ASM-NEXT: .end_amdhsa_kernel172 173// Test that we are including special SGPR usage in the granulated count.174.p2align 6175.amdhsa_kernel special_sgpr176 .amdhsa_next_free_sgpr 27177 178 .amdhsa_reserve_vcc 0179 180 .amdhsa_float_denorm_mode_16_64 0181 .amdhsa_next_free_vgpr 0182.end_amdhsa_kernel183 184// ASM: .amdhsa_kernel special_sgpr185// ASM: .amdhsa_next_free_vgpr 0186// ASM-NEXT: .amdhsa_next_free_sgpr 27187// ASM-NEXT: .amdhsa_reserve_vcc 0188// ASM: .amdhsa_float_denorm_mode_16_64 0189// ASM: .end_amdhsa_kernel190 191// Test that explicitly disabling user_sgpr's does not affect the user_sgpr192// count, i.e. this should produce the same descriptor as minimal.193.p2align 6194.amdhsa_kernel disabled_user_sgpr195 .amdhsa_next_free_vgpr 0196 .amdhsa_next_free_sgpr 0197.end_amdhsa_kernel198 199// ASM: .amdhsa_kernel disabled_user_sgpr200// ASM: .amdhsa_next_free_vgpr 0201// ASM-NEXT: .amdhsa_next_free_sgpr 0202// ASM: .end_amdhsa_kernel203 204.section .foo205 206.byte .amdgcn.gfx_generation_number207// ASM: .byte 12208 209.byte .amdgcn.gfx_generation_minor210// ASM: .byte 0211 212.byte .amdgcn.gfx_generation_stepping213// ASM: .byte 0214 215.byte .amdgcn.next_free_vgpr216// ASM: .byte 0217.byte .amdgcn.next_free_sgpr218// ASM: .byte 0219 220v_mov_b32_e32 v7, s10221 222.byte .amdgcn.next_free_vgpr223// ASM: .byte 8224.byte .amdgcn.next_free_sgpr225// ASM: .byte 11226 227.set .amdgcn.next_free_vgpr, 0228.set .amdgcn.next_free_sgpr, 0229 230.byte .amdgcn.next_free_vgpr231// ASM: .byte 0232.byte .amdgcn.next_free_sgpr233// ASM: .byte 0234 235v_mov_b32_e32 v16, s3236 237.byte .amdgcn.next_free_vgpr238// ASM: .byte 17239.byte .amdgcn.next_free_sgpr240// ASM: .byte 4241 242// Metadata243 244.amdgpu_metadata245 amdhsa.version:246 - 3247 - 0248 amdhsa.kernels:249 - .name: amd_kernel_code_t_test_all250 .symbol: amd_kernel_code_t_test_all@kd251 .kernarg_segment_size: 8252 .group_segment_fixed_size: 16253 .private_segment_fixed_size: 32254 .kernarg_segment_align: 64255 .wavefront_size: 128256 .sgpr_count: 14257 .vgpr_count: 40258 .max_flat_workgroup_size: 256259 - .name: amd_kernel_code_t_minimal260 .symbol: amd_kernel_code_t_minimal@kd261 .kernarg_segment_size: 8262 .group_segment_fixed_size: 16263 .private_segment_fixed_size: 32264 .kernarg_segment_align: 64265 .wavefront_size: 128266 .sgpr_count: 14267 .vgpr_count: 40268 .max_flat_workgroup_size: 256269.end_amdgpu_metadata270 271// ASM: .amdgpu_metadata272// ASM: amdhsa.kernels:273// ASM: - .group_segment_fixed_size: 16274// ASM: .kernarg_segment_align: 64275// ASM: .kernarg_segment_size: 8276// ASM: .max_flat_workgroup_size: 256277// ASM: .name: amd_kernel_code_t_test_all278// ASM: .private_segment_fixed_size: 32279// ASM: .sgpr_count: 14280// ASM: .symbol: 'amd_kernel_code_t_test_all@kd'281// ASM: .vgpr_count: 40282// ASM: .wavefront_size: 128283// ASM: - .group_segment_fixed_size: 16284// ASM: .kernarg_segment_align: 64285// ASM: .kernarg_segment_size: 8286// ASM: .max_flat_workgroup_size: 256287// ASM: .name: amd_kernel_code_t_minimal288// ASM: .private_segment_fixed_size: 32289// ASM: .sgpr_count: 14290// ASM: .symbol: 'amd_kernel_code_t_minimal@kd'291// ASM: .vgpr_count: 40292// ASM: .wavefront_size: 128293// ASM: amdhsa.version:294// ASM-NEXT: - 3295// ASM-NEXT: - 0296// ASM: .end_amdgpu_metadata297