318 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// RUN: llvm-mc -triple amdgcn-amd-amdhsa -mcpu=gfx904 --amdhsa-code-object-version=6 -mattr=+xnack < %s | FileCheck --check-prefix=ASM %s7// RUN: llvm-mc -triple amdgcn-amd-amdhsa -mcpu=gfx904 --amdhsa-code-object-version=6 -mattr=+xnack -filetype=obj < %s > %t8// RUN: llvm-readelf -S -r -s %t | FileCheck --check-prefix=READOBJ %s9// RUN: llvm-objdump -s -j .rodata %t | FileCheck --check-prefix=OBJDUMP %s10 11// READOBJ: Section Headers12// READOBJ: .text PROGBITS {{[0-9a-f]+}} {{[0-9a-f]+}} {{[0-9a-f]+}} {{[0-9]+}} AX {{[0-9]+}} {{[0-9]+}} 25613// READOBJ: .rodata PROGBITS {{[0-9a-f]+}} {{[0-9a-f]+}} 000100 {{[0-9]+}} A {{[0-9]+}} {{[0-9]+}} 6414 15// READOBJ: Relocation section '.rela.rodata' at offset16// READOBJ: 0000000000000010 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 1017// READOBJ: 0000000000000050 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 11018// READOBJ: 0000000000000090 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 21019// READOBJ: 00000000000000d0 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 31020 21// READOBJ: Symbol table '.symtab' contains {{[0-9]+}} entries:22// READOBJ: 0000000000000000 0 FUNC LOCAL PROTECTED 2 minimal23// READOBJ-NEXT: 0000000000000100 0 FUNC LOCAL PROTECTED 2 complete24// READOBJ-NEXT: 0000000000000200 0 FUNC LOCAL PROTECTED 2 special_sgpr25// READOBJ-NEXT: 0000000000000300 0 FUNC LOCAL PROTECTED 2 disabled_user_sgpr26// READOBJ-NEXT: 0000000000000000 64 OBJECT LOCAL DEFAULT 3 minimal.kd27// READOBJ-NEXT: 0000000000000040 64 OBJECT LOCAL DEFAULT 3 complete.kd28// READOBJ-NEXT: 0000000000000080 64 OBJECT LOCAL DEFAULT 3 special_sgpr.kd29// READOBJ-NEXT: 00000000000000c0 64 OBJECT LOCAL DEFAULT 3 disabled_user_sgpr.kd30 31// OBJDUMP: Contents of section .rodata32// Note, relocation for KERNEL_CODE_ENTRY_BYTE_OFFSET is not resolved here.33// minimal34// OBJDUMP-NEXT: 0000 00000000 00000000 00000000 0000000035// OBJDUMP-NEXT: 0010 00000000 00000000 00000000 0000000036// OBJDUMP-NEXT: 0020 00000000 00000000 00000000 0000000037// OBJDUMP-NEXT: 0030 0000ac00 80000000 00000000 0000000038// complete39// OBJDUMP-NEXT: 0040 01000000 01000000 08000000 0000000040// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 0000000041// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 0000000042// OBJDUMP-NEXT: 0070 c2500104 1f0f007f 7f080000 0000000043// special_sgpr44// OBJDUMP-NEXT: 0080 00000000 00000000 00000000 0000000045// OBJDUMP-NEXT: 0090 00000000 00000000 00000000 0000000046// OBJDUMP-NEXT: 00a0 00000000 00000000 00000000 0000000047// OBJDUMP-NEXT: 00b0 00010000 80000000 00000000 0000000048// disabled_user_sgpr49// OBJDUMP-NEXT: 00c0 00000000 00000000 00000000 0000000050// OBJDUMP-NEXT: 00d0 00000000 00000000 00000000 0000000051// OBJDUMP-NEXT: 00e0 00000000 00000000 00000000 0000000052// OBJDUMP-NEXT: 00f0 0000ac00 80000000 00000000 0000000053 54.amdgcn_target "amdgcn-amd-amdhsa--gfx904:xnack+"55// ASM: .amdgcn_target "amdgcn-amd-amdhsa--gfx904:xnack+"56 57.amdhsa_code_object_version 558// ASM: .amdhsa_code_object_version 559 60.p2align 861.type minimal,@function62minimal:63 s_endpgm64 65.p2align 866.type complete,@function67complete:68 s_endpgm69 70.p2align 871.type special_sgpr,@function72special_sgpr:73 s_endpgm74 75.p2align 876.type disabled_user_sgpr,@function77disabled_user_sgpr:78 s_endpgm79 80.rodata81// ASM: .rodata82 83// Test that only specifying required directives is allowed, and that defaulted84// values are omitted.85.p2align 686.amdhsa_kernel minimal87 .amdhsa_next_free_vgpr 088 .amdhsa_next_free_sgpr 089.end_amdhsa_kernel90 91// ASM: .amdhsa_kernel minimal92// ASM: .amdhsa_next_free_vgpr 093// ASM-NEXT: .amdhsa_next_free_sgpr 094// ASM: .end_amdhsa_kernel95 96// Test that we can specify all available directives with non-default values.97.p2align 698.amdhsa_kernel complete99 .amdhsa_group_segment_fixed_size 1100 .amdhsa_private_segment_fixed_size 1101 .amdhsa_kernarg_size 8102 .amdhsa_user_sgpr_count 15103 .amdhsa_user_sgpr_private_segment_buffer 1104 .amdhsa_user_sgpr_dispatch_ptr 1105 .amdhsa_user_sgpr_queue_ptr 1106 .amdhsa_user_sgpr_kernarg_segment_ptr 1107 .amdhsa_user_sgpr_dispatch_id 1108 .amdhsa_user_sgpr_flat_scratch_init 1109 .amdhsa_user_sgpr_private_segment_size 1110 .amdhsa_uses_dynamic_stack 1111 .amdhsa_system_sgpr_private_segment_wavefront_offset 1112 .amdhsa_system_sgpr_workgroup_id_x 0113 .amdhsa_system_sgpr_workgroup_id_y 1114 .amdhsa_system_sgpr_workgroup_id_z 1115 .amdhsa_system_sgpr_workgroup_info 1116 .amdhsa_system_vgpr_workitem_id 1117 .amdhsa_next_free_vgpr 9118 .amdhsa_next_free_sgpr 27119 .amdhsa_reserve_vcc 0120 .amdhsa_reserve_flat_scratch 0121 .amdhsa_reserve_xnack_mask 1122 .amdhsa_float_round_mode_32 1123 .amdhsa_float_round_mode_16_64 1124 .amdhsa_float_denorm_mode_32 1125 .amdhsa_float_denorm_mode_16_64 0126 .amdhsa_dx10_clamp 0127 .amdhsa_ieee_mode 0128 .amdhsa_fp16_overflow 1129 .amdhsa_exception_fp_ieee_invalid_op 1130 .amdhsa_exception_fp_denorm_src 1131 .amdhsa_exception_fp_ieee_div_zero 1132 .amdhsa_exception_fp_ieee_overflow 1133 .amdhsa_exception_fp_ieee_underflow 1134 .amdhsa_exception_fp_ieee_inexact 1135 .amdhsa_exception_int_div_zero 1136.end_amdhsa_kernel137 138// ASM: .amdhsa_kernel complete139// ASM-NEXT: .amdhsa_group_segment_fixed_size 1140// ASM-NEXT: .amdhsa_private_segment_fixed_size 1141// ASM-NEXT: .amdhsa_kernarg_size 8142// ASM-NEXT: .amdhsa_user_sgpr_count 15143// ASM-NEXT: .amdhsa_user_sgpr_private_segment_buffer 1144// ASM-NEXT: .amdhsa_user_sgpr_dispatch_ptr 1145// ASM-NEXT: .amdhsa_user_sgpr_queue_ptr 1146// ASM-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 1147// ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1148// ASM-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1149// ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1150// ASM-NEXT: .amdhsa_uses_dynamic_stack 1151// ASM-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1152// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_x 0153// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_y 1154// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z 1155// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info 1156// ASM-NEXT: .amdhsa_system_vgpr_workitem_id 1157// ASM-NEXT: .amdhsa_next_free_vgpr 9158// ASM-NEXT: .amdhsa_next_free_sgpr 27159// ASM-NEXT: .amdhsa_reserve_vcc 0160// ASM-NEXT: .amdhsa_reserve_flat_scratch 0161// ASM-NEXT: .amdhsa_reserve_xnack_mask 1162// ASM-NEXT: .amdhsa_float_round_mode_32 1163// ASM-NEXT: .amdhsa_float_round_mode_16_64 1164// ASM-NEXT: .amdhsa_float_denorm_mode_32 1165// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 0166// ASM-NEXT: .amdhsa_dx10_clamp 0167// ASM-NEXT: .amdhsa_ieee_mode 0168// ASM-NEXT: .amdhsa_fp16_overflow 1169// ASM-NEXT: .amdhsa_exception_fp_ieee_invalid_op 1170// ASM-NEXT: .amdhsa_exception_fp_denorm_src 1171// ASM-NEXT: .amdhsa_exception_fp_ieee_div_zero 1172// ASM-NEXT: .amdhsa_exception_fp_ieee_overflow 1173// ASM-NEXT: .amdhsa_exception_fp_ieee_underflow 1174// ASM-NEXT: .amdhsa_exception_fp_ieee_inexact 1175// ASM-NEXT: .amdhsa_exception_int_div_zero 1176// ASM-NEXT: .end_amdhsa_kernel177 178// Test that we are including special SGPR usage in the granulated count.179.p2align 6180.amdhsa_kernel special_sgpr181 // Same next_free_sgpr as "complete", but...182 .amdhsa_next_free_sgpr 27183 // ...on GFX9 this should require an additional 6 SGPRs, pushing us from184 // 3 granules to 4185 .amdhsa_reserve_flat_scratch 1186 187 .amdhsa_reserve_vcc 0188 .amdhsa_reserve_xnack_mask 1189 190 .amdhsa_float_denorm_mode_16_64 0191 .amdhsa_dx10_clamp 0192 .amdhsa_ieee_mode 0193 .amdhsa_next_free_vgpr 0194.end_amdhsa_kernel195 196// ASM: .amdhsa_kernel special_sgpr197// ASM: .amdhsa_next_free_vgpr 0198// ASM-NEXT: .amdhsa_next_free_sgpr 27199// ASM-NEXT: .amdhsa_reserve_vcc 0200// ASM-NEXT: .amdhsa_reserve_flat_scratch 1201// ASM-NEXT: .amdhsa_reserve_xnack_mask 1202// ASM: .amdhsa_float_denorm_mode_16_64 0203// ASM-NEXT: .amdhsa_dx10_clamp 0204// ASM-NEXT: .amdhsa_ieee_mode 0205// ASM: .end_amdhsa_kernel206 207// Test that explicitly disabling user_sgpr's does not affect the user_sgpr208// count, i.e. this should produce the same descriptor as minimal.209.p2align 6210.amdhsa_kernel disabled_user_sgpr211 .amdhsa_user_sgpr_private_segment_buffer 0212 .amdhsa_next_free_vgpr 0213 .amdhsa_next_free_sgpr 0214.end_amdhsa_kernel215 216// ASM: .amdhsa_kernel disabled_user_sgpr217// ASM: .amdhsa_next_free_vgpr 0218// ASM-NEXT: .amdhsa_next_free_sgpr 0219// ASM: .end_amdhsa_kernel220 221.section .foo222 223.byte .amdgcn.gfx_generation_number224// ASM: .byte 9225 226.byte .amdgcn.gfx_generation_minor227// ASM: .byte 0228 229.byte .amdgcn.gfx_generation_stepping230// ASM: .byte 4231 232.byte .amdgcn.next_free_vgpr233// ASM: .byte 0234.byte .amdgcn.next_free_sgpr235// ASM: .byte 0236 237v_mov_b32_e32 v7, s10238 239.byte .amdgcn.next_free_vgpr240// ASM: .byte 8241.byte .amdgcn.next_free_sgpr242// ASM: .byte 11243 244.set .amdgcn.next_free_vgpr, 0245.set .amdgcn.next_free_sgpr, 0246 247.byte .amdgcn.next_free_vgpr248// ASM: .byte 0249.byte .amdgcn.next_free_sgpr250// ASM: .byte 0251 252v_mov_b32_e32 v16, s3253 254.byte .amdgcn.next_free_vgpr255// ASM: .byte 17256.byte .amdgcn.next_free_sgpr257// ASM: .byte 4258 259// Metadata260 261.amdgpu_metadata262 amdhsa.version:263 - 3264 - 0265 amdhsa.kernels:266 - .name: amd_kernel_code_t_test_all267 .symbol: amd_kernel_code_t_test_all@kd268 .kernarg_segment_size: 8269 .group_segment_fixed_size: 16270 .private_segment_fixed_size: 32271 .uses_dynamic_stack: true272 .kernarg_segment_align: 64273 .wavefront_size: 128274 .sgpr_count: 14275 .vgpr_count: 40276 .max_flat_workgroup_size: 256277 - .name: amd_kernel_code_t_minimal278 .symbol: amd_kernel_code_t_minimal@kd279 .kernarg_segment_size: 8280 .group_segment_fixed_size: 16281 .private_segment_fixed_size: 32282 .uses_dynamic_stack: true283 .kernarg_segment_align: 64284 .wavefront_size: 128285 .sgpr_count: 14286 .vgpr_count: 40287 .max_flat_workgroup_size: 256288.end_amdgpu_metadata289 290// ASM: .amdgpu_metadata291// ASM: amdhsa.kernels:292// ASM: - .group_segment_fixed_size: 16293// ASM: .kernarg_segment_align: 64294// ASM: .kernarg_segment_size: 8295// ASM: .max_flat_workgroup_size: 256296// ASM: .name: amd_kernel_code_t_test_all297// ASM: .private_segment_fixed_size: 32298// ASM: .sgpr_count: 14299// ASM: .symbol: 'amd_kernel_code_t_test_all@kd'300// ASM: .uses_dynamic_stack: true301// ASM: .vgpr_count: 40302// ASM: .wavefront_size: 128303// ASM: - .group_segment_fixed_size: 16304// ASM: .kernarg_segment_align: 64305// ASM: .kernarg_segment_size: 8306// ASM: .max_flat_workgroup_size: 256307// ASM: .name: amd_kernel_code_t_minimal308// ASM: .private_segment_fixed_size: 32309// ASM: .sgpr_count: 14310// ASM: .symbol: 'amd_kernel_code_t_minimal@kd'311// ASM: .uses_dynamic_stack: true312// ASM: .vgpr_count: 40313// ASM: .wavefront_size: 128314// ASM: amdhsa.version:315// ASM-NEXT: - 3316// ASM-NEXT: - 0317// ASM: .end_amdgpu_metadata318