[llvm] 3d9f011 - [AMDGPU] Make the uses_dynamic_stack field in the kernel descriptor and the metadata map specific to code object v5 and later
Abinav Puthan Purayil via llvm-commits
llvm-commits at lists.llvm.org
Tue Oct 11 10:59:19 PDT 2022
Author: Abinav Puthan Purayil
Date: 2022-10-11T23:28:43+05:30
New Revision: 3d9f011a9c624b3128bc6b5e71ed456eb92e268c
URL: https://github.com/llvm/llvm-project/commit/3d9f011a9c624b3128bc6b5e71ed456eb92e268c
DIFF: https://github.com/llvm/llvm-project/commit/3d9f011a9c624b3128bc6b5e71ed456eb92e268c.diff
LOG: [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
Added:
llvm/test/MC/AMDGPU/hsa-v5-uses-dynamic-stack.s
Modified:
llvm/docs/AMDGPUUsage.rst
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
llvm/test/CodeGen/AMDGPU/gfx11-user-sgpr-init16-bug.ll
llvm/test/CodeGen/AMDGPU/recursion.ll
llvm/test/CodeGen/AMDGPU/resource-usage-dead-function.ll
llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll
llvm/test/MC/AMDGPU/hsa-gfx10-v3.s
llvm/test/MC/AMDGPU/hsa-gfx11-v3.s
llvm/test/MC/AMDGPU/hsa-gfx90a-v3.s
llvm/test/MC/AMDGPU/hsa-gfx940-v3.s
llvm/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s
llvm/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s
llvm/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s
llvm/test/MC/AMDGPU/hsa-v3.s
llvm/test/MC/AMDGPU/hsa-v4.s
llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s
Removed:
################################################################################
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 1e1765bd62627..da8a5cf19dbc9 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -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.
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
index 7e257d213a5fe..be08b7f721051 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -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;
}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp b/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
index 43a9e54de9271..2d210ecc051f8 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
@@ -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"] =
diff --git a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
index 4745a35ebdea8..43a1dfc7f561b 100644
--- a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
+++ b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
@@ -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;
diff --git a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
index 0e71509cf2bd2..1c586a68a9288 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
@@ -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"
diff --git a/llvm/test/CodeGen/AMDGPU/gfx11-user-sgpr-init16-bug.ll b/llvm/test/CodeGen/AMDGPU/gfx11-user-sgpr-init16-bug.ll
index 377a0790a184d..e48d4e5a0ee3d 100644
--- a/llvm/test/CodeGen/AMDGPU/gfx11-user-sgpr-init16-bug.ll
+++ b/llvm/test/CodeGen/AMDGPU/gfx11-user-sgpr-init16-bug.ll
@@ -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
diff --git a/llvm/test/CodeGen/AMDGPU/recursion.ll b/llvm/test/CodeGen/AMDGPU/recursion.ll
index 8e8657d0b9af9..9aedfad6fe320 100644
--- a/llvm/test/CodeGen/AMDGPU/recursion.ll
+++ b/llvm/test/CodeGen/AMDGPU/recursion.ll
@@ -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
diff --git a/llvm/test/CodeGen/AMDGPU/resource-usage-dead-function.ll b/llvm/test/CodeGen/AMDGPU/resource-usage-dead-function.ll
index f1b4b0383e3a7..1c8b8be33b0ad 100644
--- a/llvm/test/CodeGen/AMDGPU/resource-usage-dead-function.ll
+++ b/llvm/test/CodeGen/AMDGPU/resource-usage-dead-function.ll
@@ -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
diff --git a/llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll b/llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll
index a83a2e759a5af..7cc2b8214a36d 100644
--- a/llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll
+++ b/llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll
@@ -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
diff --git a/llvm/test/MC/AMDGPU/hsa-gfx10-v3.s b/llvm/test/MC/AMDGPU/hsa-gfx10-v3.s
index 32281c97c2631..ba60000837cdc 100644
--- a/llvm/test/MC/AMDGPU/hsa-gfx10-v3.s
+++ b/llvm/test/MC/AMDGPU/hsa-gfx10-v3.s
@@ -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
diff --git a/llvm/test/MC/AMDGPU/hsa-gfx11-v3.s b/llvm/test/MC/AMDGPU/hsa-gfx11-v3.s
index ca0783f0ede02..7f885b457aa63 100644
--- a/llvm/test/MC/AMDGPU/hsa-gfx11-v3.s
+++ b/llvm/test/MC/AMDGPU/hsa-gfx11-v3.s
@@ -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
diff --git a/llvm/test/MC/AMDGPU/hsa-gfx90a-v3.s b/llvm/test/MC/AMDGPU/hsa-gfx90a-v3.s
index 4c107630ee6e9..6ab70e5c72b2f 100644
--- a/llvm/test/MC/AMDGPU/hsa-gfx90a-v3.s
+++ b/llvm/test/MC/AMDGPU/hsa-gfx90a-v3.s
@@ -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
diff --git a/llvm/test/MC/AMDGPU/hsa-gfx940-v3.s b/llvm/test/MC/AMDGPU/hsa-gfx940-v3.s
index 46a730330dd13..bf2e6081a8050 100644
--- a/llvm/test/MC/AMDGPU/hsa-gfx940-v3.s
+++ b/llvm/test/MC/AMDGPU/hsa-gfx940-v3.s
@@ -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
diff --git a/llvm/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s b/llvm/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s
index 771f37dd5b85c..7da24a9d42ad4 100644
--- a/llvm/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s
+++ b/llvm/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s
@@ -39,7 +39,6 @@
; CHECK-NEXT: .private_segment_fixed_size: 32
; CHECK-NEXT: .sgpr_count: 14
; CHECK-NEXT: .symbol: 'test_kernel at 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
diff --git a/llvm/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s b/llvm/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s
index 6cdd21557cfcb..602b4566d7f25 100644
--- a/llvm/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s
+++ b/llvm/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s
@@ -20,7 +20,6 @@
// CHECK-NEXT: - 4
// CHECK: .sgpr_count: 14
// CHECK: .symbol: 'test_kernel at 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
diff --git a/llvm/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s b/llvm/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s
index 064ccb3a75f9e..cad60ea0d6d4a 100644
--- a/llvm/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s
+++ b/llvm/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s
@@ -13,7 +13,6 @@
// CHECK: .sgpr_count: 40
// CHECK: .sgpr_spill_count: 1
// CHECK: .symbol: 'test_kernel at 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
diff --git a/llvm/test/MC/AMDGPU/hsa-v3.s b/llvm/test/MC/AMDGPU/hsa-v3.s
index d9dbd5988d2e8..9f854986d7bc4 100644
--- a/llvm/test/MC/AMDGPU/hsa-v3.s
+++ b/llvm/test/MC/AMDGPU/hsa-v3.s
@@ -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 at 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 at kd'
-// ASM: .uses_dynamic_stack: true
// ASM: .vgpr_count: 40
// ASM: .wavefront_size: 128
// ASM: amdhsa.version:
diff --git a/llvm/test/MC/AMDGPU/hsa-v4.s b/llvm/test/MC/AMDGPU/hsa-v4.s
index 79dd9a06b0f03..6a824b8bcc7b9 100644
--- a/llvm/test/MC/AMDGPU/hsa-v4.s
+++ b/llvm/test/MC/AMDGPU/hsa-v4.s
@@ -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 at 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 at 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 at kd'
-// ASM: .uses_dynamic_stack: true
// ASM: .vgpr_count: 40
// ASM: .wavefront_size: 128
// ASM: amdhsa.version:
diff --git a/llvm/test/MC/AMDGPU/hsa-v5-uses-dynamic-stack.s b/llvm/test/MC/AMDGPU/hsa-v5-uses-dynamic-stack.s
new file mode 100644
index 0000000000000..6edac771faa07
--- /dev/null
+++ b/llvm/test/MC/AMDGPU/hsa-v5-uses-dynamic-stack.s
@@ -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, at function
+minimal:
+ s_endpgm
+
+.p2align 8
+.type complete, at function
+complete:
+ s_endpgm
+
+.p2align 8
+.type special_sgpr, at function
+special_sgpr:
+ s_endpgm
+
+.p2align 8
+.type disabled_user_sgpr, at 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 at 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 at 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 at 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 at 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
diff --git a/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s b/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s
index a84b2abcb1f47..73022c3c89f78 100644
--- a/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s
+++ b/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s
@@ -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 at 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 at 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
More information about the llvm-commits
mailing list