307 lines · plain
1// RUN: llvm-mc -triple amdgcn-amd-amdhsa -mcpu=gfx904 -mattr=+xnack < %s | FileCheck --check-prefix=ASM %s2// RUN: llvm-mc -triple amdgcn-amd-amdhsa -mcpu=gfx904 -mattr=+xnack -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 0000ac00 80000000 00000000 0000000033// complete34// OBJDUMP-NEXT: 0040 01000000 01000000 08000000 0000000035// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 0000000036// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 0000000037// OBJDUMP-NEXT: 0070 c2500104 1f0f007f 7f000000 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 00010000 80000000 00000000 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 0000ac00 80000000 00000000 0000000048 49.amdgcn_target "amdgcn-amd-amdhsa--gfx904:xnack+"50// ASM: .amdgcn_target "amdgcn-amd-amdhsa--gfx904:xnack+"51 52.amdhsa_code_object_version 453// ASM: .amdhsa_code_object_version 454 55.p2align 856.type minimal,@function57minimal:58 s_endpgm59 60.p2align 861.type complete,@function62complete:63 s_endpgm64 65.p2align 866.type special_sgpr,@function67special_sgpr:68 s_endpgm69 70.p2align 871.type disabled_user_sgpr,@function72disabled_user_sgpr:73 s_endpgm74 75.rodata76// ASM: .rodata77 78// Test that only specifying required directives is allowed, and that defaulted79// values are omitted.80.p2align 681.amdhsa_kernel minimal82 .amdhsa_next_free_vgpr 083 .amdhsa_next_free_sgpr 084.end_amdhsa_kernel85 86// ASM: .amdhsa_kernel minimal87// ASM: .amdhsa_next_free_vgpr 088// ASM-NEXT: .amdhsa_next_free_sgpr 089// ASM: .end_amdhsa_kernel90 91// Test that we can specify all available directives with non-default values.92.p2align 693.amdhsa_kernel complete94 .amdhsa_group_segment_fixed_size 195 .amdhsa_private_segment_fixed_size 196 .amdhsa_kernarg_size 897 .amdhsa_user_sgpr_count 1598 .amdhsa_user_sgpr_private_segment_buffer 199 .amdhsa_user_sgpr_dispatch_ptr 1100 .amdhsa_user_sgpr_queue_ptr 1101 .amdhsa_user_sgpr_kernarg_segment_ptr 1102 .amdhsa_user_sgpr_dispatch_id 1103 .amdhsa_user_sgpr_flat_scratch_init 1104 .amdhsa_user_sgpr_private_segment_size 1105 .amdhsa_system_sgpr_private_segment_wavefront_offset 1106 .amdhsa_system_sgpr_workgroup_id_x 0107 .amdhsa_system_sgpr_workgroup_id_y 1108 .amdhsa_system_sgpr_workgroup_id_z 1109 .amdhsa_system_sgpr_workgroup_info 1110 .amdhsa_system_vgpr_workitem_id 1111 .amdhsa_next_free_vgpr 9112 .amdhsa_next_free_sgpr 27113 .amdhsa_reserve_vcc 0114 .amdhsa_reserve_flat_scratch 0115 .amdhsa_reserve_xnack_mask 1116 .amdhsa_float_round_mode_32 1117 .amdhsa_float_round_mode_16_64 1118 .amdhsa_float_denorm_mode_32 1119 .amdhsa_float_denorm_mode_16_64 0120 .amdhsa_dx10_clamp 0121 .amdhsa_ieee_mode 0122 .amdhsa_fp16_overflow 1123 .amdhsa_exception_fp_ieee_invalid_op 1124 .amdhsa_exception_fp_denorm_src 1125 .amdhsa_exception_fp_ieee_div_zero 1126 .amdhsa_exception_fp_ieee_overflow 1127 .amdhsa_exception_fp_ieee_underflow 1128 .amdhsa_exception_fp_ieee_inexact 1129 .amdhsa_exception_int_div_zero 1130.end_amdhsa_kernel131 132// ASM: .amdhsa_kernel complete133// ASM-NEXT: .amdhsa_group_segment_fixed_size 1134// ASM-NEXT: .amdhsa_private_segment_fixed_size 1135// ASM-NEXT: .amdhsa_kernarg_size 8136// ASM-NEXT: .amdhsa_user_sgpr_count 15137// ASM-NEXT: .amdhsa_user_sgpr_private_segment_buffer 1138// ASM-NEXT: .amdhsa_user_sgpr_dispatch_ptr 1139// ASM-NEXT: .amdhsa_user_sgpr_queue_ptr 1140// ASM-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 1141// ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1142// ASM-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1143// ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1144// ASM-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1145// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_x 0146// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_y 1147// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z 1148// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info 1149// ASM-NEXT: .amdhsa_system_vgpr_workitem_id 1150// ASM-NEXT: .amdhsa_next_free_vgpr 9151// ASM-NEXT: .amdhsa_next_free_sgpr 27152// ASM-NEXT: .amdhsa_reserve_vcc 0153// ASM-NEXT: .amdhsa_reserve_flat_scratch 0154// ASM-NEXT: .amdhsa_reserve_xnack_mask 1155// ASM-NEXT: .amdhsa_float_round_mode_32 1156// ASM-NEXT: .amdhsa_float_round_mode_16_64 1157// ASM-NEXT: .amdhsa_float_denorm_mode_32 1158// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 0159// ASM-NEXT: .amdhsa_dx10_clamp 0160// ASM-NEXT: .amdhsa_ieee_mode 0161// ASM-NEXT: .amdhsa_fp16_overflow 1162// ASM-NEXT: .amdhsa_exception_fp_ieee_invalid_op 1163// ASM-NEXT: .amdhsa_exception_fp_denorm_src 1164// ASM-NEXT: .amdhsa_exception_fp_ieee_div_zero 1165// ASM-NEXT: .amdhsa_exception_fp_ieee_overflow 1166// ASM-NEXT: .amdhsa_exception_fp_ieee_underflow 1167// ASM-NEXT: .amdhsa_exception_fp_ieee_inexact 1168// ASM-NEXT: .amdhsa_exception_int_div_zero 1169// ASM-NEXT: .end_amdhsa_kernel170 171// Test that we are including special SGPR usage in the granulated count.172.p2align 6173.amdhsa_kernel special_sgpr174 // Same next_free_sgpr as "complete", but...175 .amdhsa_next_free_sgpr 27176 // ...on GFX9 this should require an additional 6 SGPRs, pushing us from177 // 3 granules to 4178 .amdhsa_reserve_flat_scratch 1179 180 .amdhsa_reserve_vcc 0181 .amdhsa_reserve_xnack_mask 1182 183 .amdhsa_float_denorm_mode_16_64 0184 .amdhsa_dx10_clamp 0185 .amdhsa_ieee_mode 0186 .amdhsa_next_free_vgpr 0187.end_amdhsa_kernel188 189// ASM: .amdhsa_kernel special_sgpr190// ASM: .amdhsa_next_free_vgpr 0191// ASM-NEXT: .amdhsa_next_free_sgpr 27192// ASM-NEXT: .amdhsa_reserve_vcc 0193// ASM-NEXT: .amdhsa_reserve_flat_scratch 1194// ASM-NEXT: .amdhsa_reserve_xnack_mask 1195// ASM: .amdhsa_float_denorm_mode_16_64 0196// ASM-NEXT: .amdhsa_dx10_clamp 0197// ASM-NEXT: .amdhsa_ieee_mode 0198// ASM: .end_amdhsa_kernel199 200// Test that explicitly disabling user_sgpr's does not affect the user_sgpr201// count, i.e. this should produce the same descriptor as minimal.202.p2align 6203.amdhsa_kernel disabled_user_sgpr204 .amdhsa_user_sgpr_private_segment_buffer 0205 .amdhsa_next_free_vgpr 0206 .amdhsa_next_free_sgpr 0207.end_amdhsa_kernel208 209// ASM: .amdhsa_kernel disabled_user_sgpr210// ASM: .amdhsa_next_free_vgpr 0211// ASM-NEXT: .amdhsa_next_free_sgpr 0212// ASM: .end_amdhsa_kernel213 214.section .foo215 216.byte .amdgcn.gfx_generation_number217// ASM: .byte 9218 219.byte .amdgcn.gfx_generation_minor220// ASM: .byte 0221 222.byte .amdgcn.gfx_generation_stepping223// ASM: .byte 4224 225.byte .amdgcn.next_free_vgpr226// ASM: .byte 0227.byte .amdgcn.next_free_sgpr228// ASM: .byte 0229 230v_mov_b32_e32 v7, s10231 232.byte .amdgcn.next_free_vgpr233// ASM: .byte 8234.byte .amdgcn.next_free_sgpr235// ASM: .byte 11236 237.set .amdgcn.next_free_vgpr, 0238.set .amdgcn.next_free_sgpr, 0239 240.byte .amdgcn.next_free_vgpr241// ASM: .byte 0242.byte .amdgcn.next_free_sgpr243// ASM: .byte 0244 245v_mov_b32_e32 v16, s3246 247.byte .amdgcn.next_free_vgpr248// ASM: .byte 17249.byte .amdgcn.next_free_sgpr250// ASM: .byte 4251 252// Metadata253 254.amdgpu_metadata255 amdhsa.version:256 - 3257 - 0258 amdhsa.kernels:259 - .name: amd_kernel_code_t_test_all260 .symbol: amd_kernel_code_t_test_all@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 - .name: amd_kernel_code_t_minimal270 .symbol: amd_kernel_code_t_minimal@kd271 .kernarg_segment_size: 8272 .group_segment_fixed_size: 16273 .private_segment_fixed_size: 32274 .kernarg_segment_align: 64275 .wavefront_size: 128276 .sgpr_count: 14277 .vgpr_count: 40278 .max_flat_workgroup_size: 256279.end_amdgpu_metadata280 281// ASM: .amdgpu_metadata282// ASM: amdhsa.kernels:283// 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_test_all288// ASM: .private_segment_fixed_size: 32289// ASM: .sgpr_count: 14290// ASM: .symbol: 'amd_kernel_code_t_test_all@kd'291// ASM: .vgpr_count: 40292// ASM: .wavefront_size: 128293// ASM: - .group_segment_fixed_size: 16294// ASM: .kernarg_segment_align: 64295// ASM: .kernarg_segment_size: 8296// ASM: .max_flat_workgroup_size: 256297// ASM: .name: amd_kernel_code_t_minimal298// ASM: .private_segment_fixed_size: 32299// ASM: .sgpr_count: 14300// ASM: .symbol: 'amd_kernel_code_t_minimal@kd'301// ASM: .vgpr_count: 40302// ASM: .wavefront_size: 128303// ASM: amdhsa.version:304// ASM-NEXT: - 3305// ASM-NEXT: - 0306// ASM: .end_amdgpu_metadata307