OSchip
/
llvm-project
mirror of https://github.com/THU-DSP-LAB/llvm-project.git
13
0
Fork 2
Code Issues Packages Projects Releases Wiki Activity

[AMDGPU] Make the uses_dynamic_stack field in the kernel descriptor and the metadata map specific to code object v5 and later 

Unfortunately, we have a broken handling of this in the runtime of rocm
5.3. The runtime is expected to handle this correctly when v5 becomes
the default.

Differential Revision: https://reviews.llvm.org/D134714
This commit is contained in:
Abinav Puthan Purayil 2022-09-27 11:06:35 +05:30
parent 2c94f75f00
commit 3d9f011a9c
20 changed files with 352 additions and 60 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

19
llvm/docs/AMDGPUUsage.rst
View File

@ -3219,9 +3219,6 @@ same *vendor-name*.
arguments in the
kernarg segment. Must
be a power of 2.
".uses_dynamic_stack" boolean Indicates if the generated
machine code is using a
dynamically sized stack.
".wavefront_size" integer Required Wavefront size. Must
be a power of 2.
".sgpr_count" integer Required Number of scalar
@ -3553,7 +3550,8 @@ Code Object V5 Metadata
Code object V5 metadata is the same as
:ref:`amdgpu-amdhsa-code-object-metadata-v4` with the changes defined in table
:ref:`amdgpu-amdhsa-code-object-metadata-map-table-v5` and table
:ref:`amdgpu-amdhsa-code-object-metadata-map-table-v5`, table
:ref:`amdgpu-amdhsa-code-object-kernel-metadata-map-table-v5` and table
:ref:`amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v5`.
.. table:: AMDHSA Code Object V5 Metadata Map Changes
@ -3568,6 +3566,17 @@ Code object V5 metadata is the same as
version. Currently 2.
================= ============== ========= =======================================
..
.. table:: AMDHSA Code Object V5 Kernel Metadata Map Additions
:name: amdgpu-amdhsa-code-object-kernel-metadata-map-table-v5
===================== ============= ========== =======================================
String Key Value Type Required? Description
===================== ============= ========== =======================================
".uses_dynamic_stack" boolean Indicates if the generated machine code
is using a dynamically sized stack.
===================== ============= ========== =======================================
..
.. table:: AMDHSA Code Object V5 Kernel Argument Metadata Map Additions and Changes
@ -4004,6 +4013,8 @@ The fields used by CP for code objects before V3 also match those specified in
459 1 bit USES_DYNAMIC_STACK Indicates if the generated
machine code is using a
dynamically sized stack.
This is only set in code
object v5 and later.
463:460 1 bit Reserved, must be 0.
464 1 bit RESERVED_464 Deprecated, must be 0.
467:465 3 bits Reserved, must be 0.

3
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
View File

@ -417,7 +417,8 @@ uint16_t AMDGPUAsmPrinter::getAmdhsaKernelCodeProperties(
amdhsa::KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32;
}
if (CurrentProgramInfo.DynamicCallStack) {
if (CurrentProgramInfo.DynamicCallStack &&
AMDGPU::getAmdhsaCodeObjectVersion() >= 5) {
KernelCodeProperties |= amdhsa::KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK;
}

5
llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
View File

@ -874,8 +874,9 @@ msgpack::MapDocNode MetadataStreamerMsgPackV3::getHSAKernelProps(
Kern.getDocument()->getNode(ProgramInfo.LDSSize);
Kern[".private_segment_fixed_size"] =
Kern.getDocument()->getNode(ProgramInfo.ScratchSize);
Kern[".uses_dynamic_stack"] =
Kern.getDocument()->getNode(ProgramInfo.DynamicCallStack);
if (AMDGPU::getAmdhsaCodeObjectVersion() >= 5)
Kern[".uses_dynamic_stack"] =
Kern.getDocument()->getNode(ProgramInfo.DynamicCallStack);
// FIXME: The metadata treats the minimum as 16?
Kern[".kernarg_segment_align"] =

5
llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
View File

@ -2089,8 +2089,9 @@ AMDGPUDisassembler::decodeKernelDescriptorDirective(
KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32);
}
PRINT_DIRECTIVE(".amdhsa_uses_dynamic_stack",
KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK);
if (AMDGPU::getAmdhsaCodeObjectVersion() >= 5)
PRINT_DIRECTIVE(".amdhsa_uses_dynamic_stack",
KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK);
if (TwoByteBuffer & KERNEL_CODE_PROPERTY_RESERVED1)
return MCDisassembler::Fail;

5
llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
View File

@ -367,8 +367,9 @@ void AMDGPUTargetAsmStreamer::EmitAmdhsaKernelDescriptor(
PRINT_FIELD(OS, ".amdhsa_wavefront_size32", KD,
kernel_code_properties,
amdhsa::KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32);
PRINT_FIELD(OS, ".amdhsa_uses_dynamic_stack", KD, kernel_code_properties,
amdhsa::KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK);
if (AMDGPU::getAmdhsaCodeObjectVersion() >= 5)
PRINT_FIELD(OS, ".amdhsa_uses_dynamic_stack", KD, kernel_code_properties,
amdhsa::KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK);
PRINT_FIELD(OS,
(hasArchitectedFlatScratch(STI)
? ".amdhsa_enable_private_segment"

4
llvm/test/CodeGen/AMDGPU/gfx11-user-sgpr-init16-bug.ll
View File

@ -36,7 +36,6 @@
; GCN-NEXT: .amdhsa_user_sgpr_dispatch_id 0
; GCN-NEXT: .amdhsa_user_sgpr_private_segment_size 0
; GCN-NEXT: .amdhsa_wavefront_size32
; GCN-NEXT: .amdhsa_uses_dynamic_stack 0
; GCN-NEXT: .amdhsa_enable_private_segment 0
; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_x 1
; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_y 0
@ -65,7 +64,6 @@ define amdgpu_kernel void @minimal_kernel_inputs() {
; GCN-NEXT: .amdhsa_user_sgpr_dispatch_id 0
; GCN-NEXT: .amdhsa_user_sgpr_private_segment_size 0
; GCN-NEXT: .amdhsa_wavefront_size32
; GCN-NEXT: .amdhsa_uses_dynamic_stack 0
; GCN-NEXT: .amdhsa_enable_private_segment 1
; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_x 1
; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_y 0
@ -98,7 +96,6 @@ define amdgpu_kernel void @minimal_kernel_inputs_with_stack() {
; GCN-NEXT: .amdhsa_user_sgpr_dispatch_id 0
; GCN-NEXT: .amdhsa_user_sgpr_private_segment_size 0
; GCN-NEXT: .amdhsa_wavefront_size32
; GCN-NEXT: .amdhsa_uses_dynamic_stack 0
; GCN-NEXT: .amdhsa_enable_private_segment 0
; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_x 1
; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_y 0
@ -145,7 +142,6 @@ define amdgpu_kernel void @queue_ptr() {
; GCN-NEXT: .amdhsa_user_sgpr_dispatch_id 1
; GCN-NEXT: .amdhsa_user_sgpr_private_segment_size 0
; GCN-NEXT: .amdhsa_wavefront_size32
; GCN-NEXT: .amdhsa_uses_dynamic_stack 0
; GCN-NEXT: .amdhsa_enable_private_segment 1
; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_x 1
; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_y 1

3
llvm/test/CodeGen/AMDGPU/recursion.ll
View File

@ -36,6 +36,7 @@ define void @tail_recursive_with_stack() {
;
; V5-LABEL: {{^}}calls_recursive:
; V5: .amdhsa_private_segment_fixed_size 0{{$}}
; V5: .amdhsa_uses_dynamic_stack 1
define amdgpu_kernel void @calls_recursive() {
call void @recursive()
ret void
@ -59,6 +60,7 @@ define amdgpu_kernel void @kernel_indirectly_calls_tail_recursive() {
;
; V5-LABEL: {{^}}kernel_calls_tail_recursive:
; V5: .amdhsa_private_segment_fixed_size 0{{$}}
; V5: .amdhsa_uses_dynamic_stack 1
define amdgpu_kernel void @kernel_calls_tail_recursive() {
call void @tail_recursive()
ret void
@ -69,6 +71,7 @@ define amdgpu_kernel void @kernel_calls_tail_recursive() {
;
; V5-LABEL: {{^}}kernel_calls_tail_recursive_with_stack:
; V5: .amdhsa_private_segment_fixed_size 8{{$}}
; V5: .amdhsa_uses_dynamic_stack 1
define amdgpu_kernel void @kernel_calls_tail_recursive_with_stack() {
call void @tail_recursive_with_stack()
ret void

4
llvm/test/CodeGen/AMDGPU/resource-usage-dead-function.ll
View File

@ -1,4 +1,5 @@
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -o - %s | FileCheck -check-prefix=GCN %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=5 -o - %s | FileCheck -check-prefix=GCN-V5 %s
; Make sure there's no assertion when trying to report the resource
; usage for a function which becomes dead during codegen.
@ -20,7 +21,8 @@ define internal fastcc void @unreachable() {
; GCN: s_endpgm
; GCN: .amdhsa_private_segment_fixed_size 0
; GCN: .amdhsa_uses_dynamic_stack 0
; GCN-NOT: .amdhsa_uses_dynamic_stack 0
; GCN-V5: .amdhsa_uses_dynamic_stack 0
define amdgpu_kernel void @entry() {
bb0:
br i1 false, label %bb1, label %bb2

6
llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll
View File

@ -26,7 +26,6 @@ define amdgpu_kernel void @max_alignment_128() #0 {
; VI-NEXT: .amdhsa_user_sgpr_dispatch_id 0
; VI-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
; VI-NEXT: .amdhsa_user_sgpr_private_segment_size 0
; VI-NEXT: .amdhsa_uses_dynamic_stack 0
; VI-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
; VI-NEXT: .amdhsa_system_sgpr_workgroup_id_x 1
; VI-NEXT: .amdhsa_system_sgpr_workgroup_id_y 0
@ -75,7 +74,6 @@ define amdgpu_kernel void @max_alignment_128() #0 {
; GFX9-NEXT: .amdhsa_user_sgpr_dispatch_id 0
; GFX9-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
; GFX9-NEXT: .amdhsa_user_sgpr_private_segment_size 0
; GFX9-NEXT: .amdhsa_uses_dynamic_stack 0
; GFX9-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
; GFX9-NEXT: .amdhsa_system_sgpr_workgroup_id_x 1
; GFX9-NEXT: .amdhsa_system_sgpr_workgroup_id_y 0
@ -131,7 +129,6 @@ define amdgpu_kernel void @stackrealign_attr() #1 {
; VI-NEXT: .amdhsa_user_sgpr_dispatch_id 0
; VI-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
; VI-NEXT: .amdhsa_user_sgpr_private_segment_size 0
; VI-NEXT: .amdhsa_uses_dynamic_stack 0
; VI-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
; VI-NEXT: .amdhsa_system_sgpr_workgroup_id_x 1
; VI-NEXT: .amdhsa_system_sgpr_workgroup_id_y 0
@ -180,7 +177,6 @@ define amdgpu_kernel void @stackrealign_attr() #1 {
; GFX9-NEXT: .amdhsa_user_sgpr_dispatch_id 0
; GFX9-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
; GFX9-NEXT: .amdhsa_user_sgpr_private_segment_size 0
; GFX9-NEXT: .amdhsa_uses_dynamic_stack 0
; GFX9-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
; GFX9-NEXT: .amdhsa_system_sgpr_workgroup_id_x 1
; GFX9-NEXT: .amdhsa_system_sgpr_workgroup_id_y 0
@ -236,7 +232,6 @@ define amdgpu_kernel void @alignstack_attr() #2 {
; VI-NEXT: .amdhsa_user_sgpr_dispatch_id 0
; VI-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
; VI-NEXT: .amdhsa_user_sgpr_private_segment_size 0
; VI-NEXT: .amdhsa_uses_dynamic_stack 0
; VI-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
; VI-NEXT: .amdhsa_system_sgpr_workgroup_id_x 1
; VI-NEXT: .amdhsa_system_sgpr_workgroup_id_y 0
@ -285,7 +280,6 @@ define amdgpu_kernel void @alignstack_attr() #2 {
; GFX9-NEXT: .amdhsa_user_sgpr_dispatch_id 0
; GFX9-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
; GFX9-NEXT: .amdhsa_user_sgpr_private_segment_size 0
; GFX9-NEXT: .amdhsa_uses_dynamic_stack 0
; GFX9-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
; GFX9-NEXT: .amdhsa_system_sgpr_workgroup_id_x 1
; GFX9-NEXT: .amdhsa_system_sgpr_workgroup_id_y 0

4
llvm/test/MC/AMDGPU/hsa-gfx10-v3.s
View File

@ -31,7 +31,7 @@
// OBJDUMP-NEXT: 0040 01000000 01000000 08000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0070 015001e4 1f0f007f 7f0c0000 00000000
// OBJDUMP-NEXT: 0070 015001e4 1f0f007f 7f040000 00000000
// special_sgpr
// OBJDUMP-NEXT: 0080 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0090 00000000 00000000 00000000 00000000
@ -91,7 +91,6 @@ special_sgpr:
.amdhsa_user_sgpr_flat_scratch_init 1
.amdhsa_user_sgpr_private_segment_size 1
.amdhsa_wavefront_size32 1
.amdhsa_uses_dynamic_stack 1
.amdhsa_system_sgpr_private_segment_wavefront_offset 1
.amdhsa_system_sgpr_workgroup_id_x 0
.amdhsa_system_sgpr_workgroup_id_y 1
@ -135,7 +134,6 @@ special_sgpr:
// ASM-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
// ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1
// ASM-NEXT: .amdhsa_wavefront_size32 1
// ASM-NEXT: .amdhsa_uses_dynamic_stack 1
// ASM-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_x 0
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_y 1

4
llvm/test/MC/AMDGPU/hsa-gfx11-v3.s
View File

@ -31,7 +31,7 @@
// OBJDUMP-NEXT: 0040 01000000 01000000 08000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0070 015001e4 130f007f 5e0c0000 00000000
// OBJDUMP-NEXT: 0070 015001e4 130f007f 5e040000 00000000
// special_sgpr
// OBJDUMP-NEXT: 0080 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0090 00000000 00000000 00000000 00000000
@ -87,7 +87,6 @@ special_sgpr:
.amdhsa_user_sgpr_dispatch_id 1
.amdhsa_user_sgpr_private_segment_size 1
.amdhsa_wavefront_size32 1
.amdhsa_uses_dynamic_stack 1
.amdhsa_enable_private_segment 1
.amdhsa_system_sgpr_workgroup_id_x 0
.amdhsa_system_sgpr_workgroup_id_y 1
@ -127,7 +126,6 @@ special_sgpr:
// ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1
// ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1
// ASM-NEXT: .amdhsa_wavefront_size32 1
// ASM-NEXT: .amdhsa_uses_dynamic_stack 1
// ASM-NEXT: .amdhsa_enable_private_segment 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_x 0
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_y 1

4
llvm/test/MC/AMDGPU/hsa-gfx90a-v3.s
View File

@ -28,7 +28,7 @@
// OBJDUMP-NEXT: 0040 01000000 01000000 00000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000100
// OBJDUMP-NEXT: 0070 c1500104 1f0f007f 7f080000 00000000
// OBJDUMP-NEXT: 0070 c1500104 1f0f007f 7f000000 00000000
.text
// ASM: .text
@ -77,7 +77,6 @@ complete:
.amdhsa_user_sgpr_dispatch_id 1
.amdhsa_user_sgpr_flat_scratch_init 1
.amdhsa_user_sgpr_private_segment_size 1
.amdhsa_uses_dynamic_stack 1
.amdhsa_system_sgpr_private_segment_wavefront_offset 1
.amdhsa_system_sgpr_workgroup_id_x 0
.amdhsa_system_sgpr_workgroup_id_y 1
@ -118,7 +117,6 @@ complete:
// ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1
// ASM-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
// ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1
// ASM-NEXT: .amdhsa_uses_dynamic_stack 1
// ASM-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_x 0
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_y 1

4
llvm/test/MC/AMDGPU/hsa-gfx940-v3.s
View File

@ -28,7 +28,7 @@
// OBJDUMP-NEXT: 0040 01000000 01000000 00000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000100
// OBJDUMP-NEXT: 0070 01510104 130f007f 5e080000 00000000
// OBJDUMP-NEXT: 0070 01510104 130f007f 5e000000 00000000
.text
// ASM: .text
@ -75,7 +75,6 @@ complete:
.amdhsa_user_sgpr_kernarg_segment_ptr 1
.amdhsa_user_sgpr_dispatch_id 1
.amdhsa_user_sgpr_private_segment_size 1
.amdhsa_uses_dynamic_stack 1
.amdhsa_enable_private_segment 1
.amdhsa_system_sgpr_workgroup_id_x 0
.amdhsa_system_sgpr_workgroup_id_y 1
@ -113,7 +112,6 @@ complete:
// ASM-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 1
// ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1
// ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1
// ASM-NEXT: .amdhsa_uses_dynamic_stack 1
// ASM-NEXT: .amdhsa_enable_private_segment 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_x 0
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_y 1

2
llvm/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s
View File

@ -39,7 +39,6 @@
; CHECK-NEXT: .private_segment_fixed_size: 32
; CHECK-NEXT: .sgpr_count: 14
; CHECK-NEXT: .symbol: 'test_kernel@kd'
; CHECK-NEXT: .uses_dynamic_stack: true
; CHECK-NEXT: .vgpr_count: 40
; CHECK-NEXT: .wavefront_size: 128
; CHECK-NEXT: amdhsa.printf:
@ -66,7 +65,6 @@
.kernarg_segment_size: 8
.group_segment_fixed_size: 16
.private_segment_fixed_size: 32
.uses_dynamic_stack: true
.kernarg_segment_align: 64
.wavefront_size: 128
.sgpr_count: 14

2
llvm/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s
View File

@ -20,7 +20,6 @@
// CHECK-NEXT: - 4
// CHECK: .sgpr_count: 14
// CHECK: .symbol: 'test_kernel@kd'
// CHECK: .uses_dynamic_stack: true
// CHECK: .vec_type_hint: int
// CHECK: .vgpr_count: 40
// CHECK: .wavefront_size: 128
@ -52,7 +51,6 @@
.kernarg_segment_size: 8
.group_segment_fixed_size: 16
.private_segment_fixed_size: 32
.uses_dynamic_stack: true
.kernarg_segment_align: 64
.wavefront_size: 128
.sgpr_count: 14

2
llvm/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s
View File

@ -13,7 +13,6 @@
// CHECK: .sgpr_count: 40
// CHECK: .sgpr_spill_count: 1
// CHECK: .symbol: 'test_kernel@kd'
// CHECK: .uses_dynamic_stack: true
// CHECK: .vgpr_count: 14
// CHECK: .vgpr_spill_count: 1
// CHECK: .wavefront_size: 64
@ -33,7 +32,6 @@
.kernarg_segment_size: 24
.group_segment_fixed_size: 24
.private_segment_fixed_size: 16
.uses_dynamic_stack: true
.kernarg_segment_align: 16
.wavefront_size: 64
.max_flat_workgroup_size: 256

8
llvm/test/MC/AMDGPU/hsa-v3.s
View File

@ -34,7 +34,7 @@
// OBJDUMP-NEXT: 0040 01000000 01000000 08000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0070 c2500104 1f0f007f 7f080000 00000000
// OBJDUMP-NEXT: 0070 c2500104 1f0f007f 7f000000 00000000
// special_sgpr
// OBJDUMP-NEXT: 0080 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0090 00000000 00000000 00000000 00000000
@ -101,7 +101,6 @@ disabled_user_sgpr:
.amdhsa_user_sgpr_dispatch_id 1
.amdhsa_user_sgpr_flat_scratch_init 1
.amdhsa_user_sgpr_private_segment_size 1
.amdhsa_uses_dynamic_stack 1
.amdhsa_system_sgpr_private_segment_wavefront_offset 1
.amdhsa_system_sgpr_workgroup_id_x 0
.amdhsa_system_sgpr_workgroup_id_y 1
@ -141,7 +140,6 @@ disabled_user_sgpr:
// ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1
// ASM-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
// ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1
// ASM-NEXT: .amdhsa_uses_dynamic_stack 1
// ASM-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_x 0
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_y 1
@ -261,7 +259,6 @@ v_mov_b32_e32 v16, s3
.kernarg_segment_size: 8
.group_segment_fixed_size: 16
.private_segment_fixed_size: 32
.uses_dynamic_stack: true
.kernarg_segment_align: 64
.wavefront_size: 128
.sgpr_count: 14
@ -272,7 +269,6 @@ v_mov_b32_e32 v16, s3
.kernarg_segment_size: 8
.group_segment_fixed_size: 16
.private_segment_fixed_size: 32
.uses_dynamic_stack: true
.kernarg_segment_align: 64
.wavefront_size: 128
.sgpr_count: 14
@ -290,7 +286,6 @@ v_mov_b32_e32 v16, s3
// ASM: .private_segment_fixed_size: 32
// ASM: .sgpr_count: 14
// ASM: .symbol: 'amd_kernel_code_t_test_all@kd'
// ASM: .uses_dynamic_stack: true
// ASM: .vgpr_count: 40
// ASM: .wavefront_size: 128
// ASM: - .group_segment_fixed_size: 16
@ -301,7 +296,6 @@ v_mov_b32_e32 v16, s3
// ASM: .private_segment_fixed_size: 32
// ASM: .sgpr_count: 14
// ASM: .symbol: 'amd_kernel_code_t_minimal@kd'
// ASM: .uses_dynamic_stack: true
// ASM: .vgpr_count: 40
// ASM: .wavefront_size: 128
// ASM: amdhsa.version:

8
llvm/test/MC/AMDGPU/hsa-v4.s
View File

@ -34,7 +34,7 @@
// OBJDUMP-NEXT: 0040 01000000 01000000 08000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0070 c2500104 1f0f007f 7f080000 00000000
// OBJDUMP-NEXT: 0070 c2500104 1f0f007f 7f000000 00000000
// special_sgpr
// OBJDUMP-NEXT: 0080 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0090 00000000 00000000 00000000 00000000
@ -102,7 +102,6 @@ disabled_user_sgpr:
.amdhsa_user_sgpr_dispatch_id 1
.amdhsa_user_sgpr_flat_scratch_init 1
.amdhsa_user_sgpr_private_segment_size 1
.amdhsa_uses_dynamic_stack 1
.amdhsa_system_sgpr_private_segment_wavefront_offset 1
.amdhsa_system_sgpr_workgroup_id_x 0
.amdhsa_system_sgpr_workgroup_id_y 1
@ -142,7 +141,6 @@ disabled_user_sgpr:
// ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1
// ASM-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
// ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1
// ASM-NEXT: .amdhsa_uses_dynamic_stack 1
// ASM-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_x 0
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_y 1
@ -262,7 +260,6 @@ v_mov_b32_e32 v16, s3
.kernarg_segment_size: 8
.group_segment_fixed_size: 16
.private_segment_fixed_size: 32
.uses_dynamic_stack: true
.kernarg_segment_align: 64
.wavefront_size: 128
.sgpr_count: 14
@ -272,7 +269,6 @@ v_mov_b32_e32 v16, s3
.symbol: amd_kernel_code_t_minimal@kd
.kernarg_segment_size: 8
.group_segment_fixed_size: 16
.uses_dynamic_stack: true
.private_segment_fixed_size: 32
.kernarg_segment_align: 64
.wavefront_size: 128
@ -291,7 +287,6 @@ v_mov_b32_e32 v16, s3
// ASM: .private_segment_fixed_size: 32
// ASM: .sgpr_count: 14
// ASM: .symbol: 'amd_kernel_code_t_test_all@kd'
// ASM: .uses_dynamic_stack: true
// ASM: .vgpr_count: 40
// ASM: .wavefront_size: 128
// ASM: - .group_segment_fixed_size: 16
@ -302,7 +297,6 @@ v_mov_b32_e32 v16, s3
// ASM: .private_segment_fixed_size: 32
// ASM: .sgpr_count: 14
// ASM: .symbol: 'amd_kernel_code_t_minimal@kd'
// ASM: .uses_dynamic_stack: true
// ASM: .vgpr_count: 40
// ASM: .wavefront_size: 128
// ASM: amdhsa.version:

311
llvm/test/MC/AMDGPU/hsa-v5-uses-dynamic-stack.s Normal file
View File

@ -0,0 +1,311 @@
// RUN: llvm-mc -triple amdgcn-amd-amdhsa -mcpu=gfx904 --amdhsa-code-object-version=5 -mattr=+xnack < %s | FileCheck --check-prefix=ASM %s
// RUN: llvm-mc -triple amdgcn-amd-amdhsa -mcpu=gfx904 --amdhsa-code-object-version=5 -mattr=+xnack -filetype=obj < %s > %t
// RUN: llvm-readelf -S -r -s %t | FileCheck --check-prefix=READOBJ %s
// RUN: llvm-objdump -s -j .rodata %t | FileCheck --check-prefix=OBJDUMP %s
// READOBJ: Section Headers
// READOBJ: .text PROGBITS {{[0-9a-f]+}} {{[0-9a-f]+}} {{[0-9a-f]+}} {{[0-9]+}} AX {{[0-9]+}} {{[0-9]+}} 256
// READOBJ: .rodata PROGBITS {{[0-9a-f]+}} {{[0-9a-f]+}} 000100 {{[0-9]+}} A {{[0-9]+}} {{[0-9]+}} 64
// READOBJ: Relocation section '.rela.rodata' at offset
// READOBJ: 0000000000000010 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 10
// READOBJ: 0000000000000050 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 110
// READOBJ: 0000000000000090 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 210
// READOBJ: 00000000000000d0 {{[0-9a-f]+}}00000005 R_AMDGPU_REL64 0000000000000000 .text + 310
// READOBJ: Symbol table '.symtab' contains {{[0-9]+}} entries:
// READOBJ: 0000000000000000 0 FUNC LOCAL PROTECTED 2 minimal
// READOBJ-NEXT: 0000000000000100 0 FUNC LOCAL PROTECTED 2 complete
// READOBJ-NEXT: 0000000000000200 0 FUNC LOCAL PROTECTED 2 special_sgpr
// READOBJ-NEXT: 0000000000000300 0 FUNC LOCAL PROTECTED 2 disabled_user_sgpr
// READOBJ-NEXT: 0000000000000000 64 OBJECT LOCAL DEFAULT 3 minimal.kd
// READOBJ-NEXT: 0000000000000040 64 OBJECT LOCAL DEFAULT 3 complete.kd
// READOBJ-NEXT: 0000000000000080 64 OBJECT LOCAL DEFAULT 3 special_sgpr.kd
// READOBJ-NEXT: 00000000000000c0 64 OBJECT LOCAL DEFAULT 3 disabled_user_sgpr.kd
// OBJDUMP: Contents of section .rodata
// Note, relocation for KERNEL_CODE_ENTRY_BYTE_OFFSET is not resolved here.
// minimal
// OBJDUMP-NEXT: 0000 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0010 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0020 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0030 0000ac00 80000000 00000000 00000000
// complete
// OBJDUMP-NEXT: 0040 01000000 01000000 08000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0070 c2500104 1f0f007f 7f080000 00000000
// special_sgpr
// OBJDUMP-NEXT: 0080 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0090 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 00a0 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 00b0 00010000 80000000 00000000 00000000
// disabled_user_sgpr
// OBJDUMP-NEXT: 00c0 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 00d0 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 00e0 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 00f0 0000ac00 80000000 00000000 00000000
.text
// ASM: .text
.amdgcn_target "amdgcn-amd-amdhsa--gfx904:xnack+"
// ASM: .amdgcn_target "amdgcn-amd-amdhsa--gfx904:xnack+"
.p2align 8
.type minimal,@function
minimal:
s_endpgm
.p2align 8
.type complete,@function
complete:
s_endpgm
.p2align 8
.type special_sgpr,@function
special_sgpr:
s_endpgm
.p2align 8
.type disabled_user_sgpr,@function
disabled_user_sgpr:
s_endpgm
.rodata
// ASM: .rodata
// Test that only specifying required directives is allowed, and that defaulted
// values are omitted.
.p2align 6
.amdhsa_kernel minimal
.amdhsa_next_free_vgpr 0
.amdhsa_next_free_sgpr 0
.end_amdhsa_kernel
// ASM: .amdhsa_kernel minimal
// ASM: .amdhsa_next_free_vgpr 0
// ASM-NEXT: .amdhsa_next_free_sgpr 0
// ASM: .end_amdhsa_kernel
// Test that we can specify all available directives with non-default values.
.p2align 6
.amdhsa_kernel complete
.amdhsa_group_segment_fixed_size 1
.amdhsa_private_segment_fixed_size 1
.amdhsa_kernarg_size 8
.amdhsa_user_sgpr_count 15
.amdhsa_user_sgpr_private_segment_buffer 1
.amdhsa_user_sgpr_dispatch_ptr 1
.amdhsa_user_sgpr_queue_ptr 1
.amdhsa_user_sgpr_kernarg_segment_ptr 1
.amdhsa_user_sgpr_dispatch_id 1
.amdhsa_user_sgpr_flat_scratch_init 1
.amdhsa_user_sgpr_private_segment_size 1
.amdhsa_uses_dynamic_stack 1
.amdhsa_system_sgpr_private_segment_wavefront_offset 1
.amdhsa_system_sgpr_workgroup_id_x 0
.amdhsa_system_sgpr_workgroup_id_y 1
.amdhsa_system_sgpr_workgroup_id_z 1
.amdhsa_system_sgpr_workgroup_info 1
.amdhsa_system_vgpr_workitem_id 1
.amdhsa_next_free_vgpr 9
.amdhsa_next_free_sgpr 27
.amdhsa_reserve_vcc 0
.amdhsa_reserve_flat_scratch 0
.amdhsa_reserve_xnack_mask 1
.amdhsa_float_round_mode_32 1
.amdhsa_float_round_mode_16_64 1
.amdhsa_float_denorm_mode_32 1
.amdhsa_float_denorm_mode_16_64 0
.amdhsa_dx10_clamp 0
.amdhsa_ieee_mode 0
.amdhsa_fp16_overflow 1
.amdhsa_exception_fp_ieee_invalid_op 1
.amdhsa_exception_fp_denorm_src 1
.amdhsa_exception_fp_ieee_div_zero 1
.amdhsa_exception_fp_ieee_overflow 1
.amdhsa_exception_fp_ieee_underflow 1
.amdhsa_exception_fp_ieee_inexact 1
.amdhsa_exception_int_div_zero 1
.end_amdhsa_kernel
// ASM: .amdhsa_kernel complete
// ASM-NEXT: .amdhsa_group_segment_fixed_size 1
// ASM-NEXT: .amdhsa_private_segment_fixed_size 1
// ASM-NEXT: .amdhsa_kernarg_size 8
// ASM-NEXT: .amdhsa_user_sgpr_count 15
// ASM-NEXT: .amdhsa_user_sgpr_private_segment_buffer 1
// ASM-NEXT: .amdhsa_user_sgpr_dispatch_ptr 1
// ASM-NEXT: .amdhsa_user_sgpr_queue_ptr 1
// ASM-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 1
// ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1
// ASM-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
// ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1
// ASM-NEXT: .amdhsa_uses_dynamic_stack 1
// ASM-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_x 0
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_y 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info 1
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id 1
// ASM-NEXT: .amdhsa_next_free_vgpr 9
// ASM-NEXT: .amdhsa_next_free_sgpr 27
// ASM-NEXT: .amdhsa_reserve_vcc 0
// ASM-NEXT: .amdhsa_reserve_flat_scratch 0
// ASM-NEXT: .amdhsa_reserve_xnack_mask 1
// ASM-NEXT: .amdhsa_float_round_mode_32 1
// ASM-NEXT: .amdhsa_float_round_mode_16_64 1
// ASM-NEXT: .amdhsa_float_denorm_mode_32 1
// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 0
// ASM-NEXT: .amdhsa_dx10_clamp 0
// ASM-NEXT: .amdhsa_ieee_mode 0
// ASM-NEXT: .amdhsa_fp16_overflow 1
// ASM-NEXT: .amdhsa_exception_fp_ieee_invalid_op 1
// ASM-NEXT: .amdhsa_exception_fp_denorm_src 1
// ASM-NEXT: .amdhsa_exception_fp_ieee_div_zero 1
// ASM-NEXT: .amdhsa_exception_fp_ieee_overflow 1
// ASM-NEXT: .amdhsa_exception_fp_ieee_underflow 1
// ASM-NEXT: .amdhsa_exception_fp_ieee_inexact 1
// ASM-NEXT: .amdhsa_exception_int_div_zero 1
// ASM-NEXT: .end_amdhsa_kernel
// Test that we are including special SGPR usage in the granulated count.
.p2align 6
.amdhsa_kernel special_sgpr
// Same next_free_sgpr as "complete", but...
.amdhsa_next_free_sgpr 27
// ...on GFX9 this should require an additional 6 SGPRs, pushing us from
// 3 granules to 4
.amdhsa_reserve_flat_scratch 1
.amdhsa_reserve_vcc 0
.amdhsa_reserve_xnack_mask 1
.amdhsa_float_denorm_mode_16_64 0
.amdhsa_dx10_clamp 0
.amdhsa_ieee_mode 0
.amdhsa_next_free_vgpr 0
.end_amdhsa_kernel
// ASM: .amdhsa_kernel special_sgpr
// ASM: .amdhsa_next_free_vgpr 0
// ASM-NEXT: .amdhsa_next_free_sgpr 27
// ASM-NEXT: .amdhsa_reserve_vcc 0
// ASM-NEXT: .amdhsa_reserve_xnack_mask 1
// ASM: .amdhsa_float_denorm_mode_16_64 0
// ASM-NEXT: .amdhsa_dx10_clamp 0
// ASM-NEXT: .amdhsa_ieee_mode 0
// ASM: .end_amdhsa_kernel
// Test that explicitly disabling user_sgpr's does not affect the user_sgpr
// count, i.e. this should produce the same descriptor as minimal.
.p2align 6
.amdhsa_kernel disabled_user_sgpr
.amdhsa_user_sgpr_private_segment_buffer 0
.amdhsa_next_free_vgpr 0
.amdhsa_next_free_sgpr 0
.end_amdhsa_kernel
// ASM: .amdhsa_kernel disabled_user_sgpr
// ASM: .amdhsa_next_free_vgpr 0
// ASM-NEXT: .amdhsa_next_free_sgpr 0
// ASM: .end_amdhsa_kernel
.section .foo
.byte .amdgcn.gfx_generation_number
// ASM: .byte 9
.byte .amdgcn.gfx_generation_minor
// ASM: .byte 0
.byte .amdgcn.gfx_generation_stepping
// ASM: .byte 4
.byte .amdgcn.next_free_vgpr
// ASM: .byte 0
.byte .amdgcn.next_free_sgpr
// ASM: .byte 0
v_mov_b32_e32 v7, s10
.byte .amdgcn.next_free_vgpr
// ASM: .byte 8
.byte .amdgcn.next_free_sgpr
// ASM: .byte 11
.set .amdgcn.next_free_vgpr, 0
.set .amdgcn.next_free_sgpr, 0
.byte .amdgcn.next_free_vgpr
// ASM: .byte 0
.byte .amdgcn.next_free_sgpr
// ASM: .byte 0
v_mov_b32_e32 v16, s3
.byte .amdgcn.next_free_vgpr
// ASM: .byte 17
.byte .amdgcn.next_free_sgpr
// ASM: .byte 4
// Metadata
.amdgpu_metadata
amdhsa.version:
- 3
- 0
amdhsa.kernels:
- .name: amd_kernel_code_t_test_all
.symbol: amd_kernel_code_t_test_all@kd
.kernarg_segment_size: 8
.group_segment_fixed_size: 16
.private_segment_fixed_size: 32
.uses_dynamic_stack: true
.kernarg_segment_align: 64
.wavefront_size: 128
.sgpr_count: 14
.vgpr_count: 40
.max_flat_workgroup_size: 256
- .name: amd_kernel_code_t_minimal
.symbol: amd_kernel_code_t_minimal@kd
.kernarg_segment_size: 8
.group_segment_fixed_size: 16
.private_segment_fixed_size: 32
.uses_dynamic_stack: true
.kernarg_segment_align: 64
.wavefront_size: 128
.sgpr_count: 14
.vgpr_count: 40
.max_flat_workgroup_size: 256
.end_amdgpu_metadata
// ASM: .amdgpu_metadata
// ASM: amdhsa.kernels:
// ASM: - .group_segment_fixed_size: 16
// ASM: .kernarg_segment_align: 64
// ASM: .kernarg_segment_size: 8
// ASM: .max_flat_workgroup_size: 256
// ASM: .name: amd_kernel_code_t_test_all
// ASM: .private_segment_fixed_size: 32
// ASM: .sgpr_count: 14
// ASM: .symbol: 'amd_kernel_code_t_test_all@kd'
// ASM: .uses_dynamic_stack: true
// ASM: .vgpr_count: 40
// ASM: .wavefront_size: 128
// ASM: - .group_segment_fixed_size: 16
// ASM: .kernarg_segment_align: 64
// ASM: .kernarg_segment_size: 8
// ASM: .max_flat_workgroup_size: 256
// ASM: .name: amd_kernel_code_t_minimal
// ASM: .private_segment_fixed_size: 32
// ASM: .sgpr_count: 14
// ASM: .symbol: 'amd_kernel_code_t_minimal@kd'
// ASM: .uses_dynamic_stack: true
// ASM: .vgpr_count: 40
// ASM: .wavefront_size: 128
// ASM: amdhsa.version:
// ASM-NEXT: - 3
// ASM-NEXT: - 0
// ASM: .end_amdgpu_metadata

9
llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s
View File

@ -10,10 +10,10 @@
#LLVM-NEXT: NoteSection {
#LLVM-NEXT: Name: .note
#LLVM-NEXT: Offset: 0x40
#LLVM-NEXT: Size: 0x128
#LLVM-NEXT: Size: 0x110
#LLVM-NEXT: Note {
#LLVM-NEXT: Owner: AMDGPU
#LLVM-NEXT: Data size: 0x111
#LLVM-NEXT: Data size: 0xFC
#LLVM-NEXT: Type: NT_AMDGPU_METADATA (AMDGPU Metadata)
#LLVM-NEXT: AMDGPU Metadata: ---
#LLVM-NEXT: amdhsa.kernels:
@ -25,7 +25,6 @@
#LLVM-NEXT: .private_segment_fixed_size: 32
#LLVM-NEXT: .sgpr_count: 14
#LLVM-NEXT: .symbol: 'test_kernel@kd'
#LLVM-NEXT: .uses_dynamic_stack: true
#LLVM-NEXT: .vgpr_count: 40
#LLVM-NEXT: .wavefront_size: 128
#LLVM-NEXT: amdhsa.version:
@ -39,7 +38,7 @@
# GNU: Displaying notes found in: .note
# GNU-NEXT: Owner Data size Description
# GNU-NEXT: AMDGPU 0x00000111 NT_AMDGPU_METADATA (AMDGPU Metadata)
# GNU-NEXT: AMDGPU 0x000000fc NT_AMDGPU_METADATA (AMDGPU Metadata)
# GNU-NEXT: AMDGPU Metadata:
# GNU-NEXT: ---
# GNU-NEXT: amdhsa.kernels:
@ -51,7 +50,6 @@
# GNU-NEXT: .private_segment_fixed_size: 32
# GNU-NEXT: .sgpr_count: 14
# GNU-NEXT: .symbol: 'test_kernel@kd'
# GNU-NEXT: .uses_dynamic_stack: true
# GNU-NEXT: .vgpr_count: 40
# GNU-NEXT: .wavefront_size: 128
# GNU-NEXT: amdhsa.version:
@ -71,7 +69,6 @@
.kernarg_segment_size: 8
.max_flat_workgroup_size: 256
.private_segment_fixed_size: 32
.uses_dynamic_stack: true
.sgpr_count: 14
.vgpr_count: 40
.wavefront_size: 128