[AMDGPU] Add the uses_dynamic_stack field to the kernel descriptor and the kernel...

This change introduces the dynamic stack boolean field to code-object-v3
and above under the code properties of the kernel descriptor and under
the kernel metadata map of NT_AMDGPU_METADATA. This field corresponds to
the is_dynamic_callstack field of amd_kernel_code_t.

Differential Revision: https://reviews.llvm.org/D128344

diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 26cba99..f591e60 100644 (file)
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -3219,6 +3219,9 @@ 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
@@ -3997,7 +4000,10 @@ The fields used by CP for code objects before V3 also match those specified in
                                                        - If 1 execute in
                                                          native wavefront size
                                                          32 mode.
-     463:459 1 bit                                   Reserved, must be 0.
+     459     1 bit   USES_DYNAMIC_STACK              Indicates if the generated
+                                                     machine code is using a
+                                                     dynamically sized stack.
+     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.
      468     1 bit   RESERVED_468                    Deprecated, must be 0.
@@ -14847,6 +14853,8 @@ terminated by an ``.end_amdhsa_kernel`` directive.
                                                               Feature                          :ref:`amdgpu-amdhsa-kernel-descriptor-v3-table`.
                                                               Specific
                                                               (wavefrontsize64)
+     ``.amdhsa_uses_dynamic_stack``                           0                   GFX6-GFX11   Controls USES_DYNAMIC_STACK in
+                                                                                               :ref:`amdgpu-amdhsa-kernel-descriptor-v3-table`.
      ``.amdhsa_system_sgpr_private_segment_wavefront_offset`` 0                   GFX6-GFX10   Controls ENABLE_PRIVATE_SEGMENT in
                                                                                   (except      :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx11-table`.
                                                                                   GFX940)

diff --git a/llvm/include/llvm/Support/AMDHSAKernelDescriptor.h b/llvm/include/llvm/Support/AMDHSAKernelDescriptor.h
index 41d144c..61b0574 100644 (file)
--- a/llvm/include/llvm/Support/AMDHSAKernelDescriptor.h
+++ b/llvm/include/llvm/Support/AMDHSAKernelDescriptor.h
@@ -161,7 +161,8 @@ enum : int32_t {
   KERNEL_CODE_PROPERTY(ENABLE_SGPR_PRIVATE_SEGMENT_SIZE, 6, 1),
   KERNEL_CODE_PROPERTY(RESERVED0, 7, 3),
   KERNEL_CODE_PROPERTY(ENABLE_WAVEFRONT_SIZE32, 10, 1), // GFX10+
-  KERNEL_CODE_PROPERTY(RESERVED1, 11, 5),
+  KERNEL_CODE_PROPERTY(USES_DYNAMIC_STACK, 11, 1),
+  KERNEL_CODE_PROPERTY(RESERVED1, 12, 4),
 };
 #undef KERNEL_CODE_PROPERTY
 

diff --git a/llvm/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp b/llvm/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp
index 1613e7e..c5ab35d 100644 (file)
--- a/llvm/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp
+++ b/llvm/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp
@@ -260,6 +260,9 @@ bool MetadataVerifier::verifyKernel(msgpack::DocNode &Node) {
     return false;
   if (!verifyIntegerEntry(KernelMap, ".private_segment_fixed_size", true))
     return false;
+  if (!verifyScalarEntry(KernelMap, ".uses_dynamic_stack", false,
+                         msgpack::Type::Boolean))
+    return false;
   if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_align", true))
     return false;
   if (!verifyIntegerEntry(KernelMap, ".wavefront_size", true))

diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
index f1cc40b..13a65f1 100644 (file)
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -417,6 +417,10 @@ uint16_t AMDGPUAsmPrinter::getAmdhsaKernelCodeProperties(
         amdhsa::KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32;
   }
 
+  if (CurrentProgramInfo.DynamicCallStack) {
+    KernelCodeProperties |= amdhsa::KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK;
+  }
+
   return KernelCodeProperties;
 }
 

diff --git a/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp b/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
index 6fa44ff..632a76b 100644 (file)
--- a/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
@@ -875,6 +875,8 @@ MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
       Kern.getDocument()->getNode(ProgramInfo.LDSSize);
   Kern[".private_segment_fixed_size"] =
       Kern.getDocument()->getNode(ProgramInfo.ScratchSize);
+  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/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index d38182b..201a7ff 100644 (file)
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -5001,6 +5001,9 @@ bool AMDGPUAsmParser::ParseDirectiveAMDHSAKernel() {
       PARSE_BITS_ENTRY(KD.kernel_code_properties,
                        KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32,
                        Val, ValRange);
+    } else if (ID == ".amdhsa_uses_dynamic_stack") {
+      PARSE_BITS_ENTRY(KD.kernel_code_properties,
+                       KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK, Val, ValRange);
     } else if (ID == ".amdhsa_system_sgpr_private_segment_wavefront_offset") {
       if (hasArchitectedFlatScratch())
         return Error(IDRange.Start,

diff --git a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
index 0f7545c..98ee720 100644 (file)
--- a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
+++ b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
@@ -2040,6 +2040,9 @@ AMDGPUDisassembler::decodeKernelDescriptorDirective(
                       KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32);
     }
 
+    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 0781334..0e71509 100644 (file)
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
@@ -367,6 +367,8 @@ 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);
   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 831d94b..178b246 100644 (file)
--- a/llvm/test/CodeGen/AMDGPU/gfx11-user-sgpr-init16-bug.ll
+++ b/llvm/test/CodeGen/AMDGPU/gfx11-user-sgpr-init16-bug.ll
@@ -21,6 +21,7 @@
 ; GCN-NEXT: .amdhsa_user_sgpr_dispatch_id 0
 ; GCN-NEXT: .amdhsa_user_sgpr_private_segment_size 0
 ; GCN-NEXT: .amdhsa_wavefront_size32 1
+; 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
@@ -46,6 +47,7 @@ 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 1
+; 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
@@ -75,6 +77,7 @@ 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 1
+; 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
@@ -115,6 +118,7 @@ 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 1
+; 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/indirect-call-known-callees.ll b/llvm/test/CodeGen/AMDGPU/indirect-call-known-callees.ll
index aebbe33..8e93f3a 100644 (file)
--- a/llvm/test/CodeGen/AMDGPU/indirect-call-known-callees.ll
+++ b/llvm/test/CodeGen/AMDGPU/indirect-call-known-callees.ll
@@ -65,6 +65,7 @@ define amdgpu_kernel void @indirect_call_known_no_special_inputs() {
 ; CHECK-NEXT: .amdhsa_user_sgpr_dispatch_id 0
 ; CHECK-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
 ; CHECK-NEXT: .amdhsa_user_sgpr_private_segment_size 0
+; CHECK-NEXT: .amdhsa_uses_dynamic_stack 1
 ; CHECK-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
 ; CHECK-NEXT: .amdhsa_system_sgpr_workgroup_id_x 1
 ; CHECK-NEXT: .amdhsa_system_sgpr_workgroup_id_y 0

diff --git a/llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll b/llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll
index 7cc2b82..a83a2e7 100644 (file)
--- a/llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll
+++ b/llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll
@@ -26,6 +26,7 @@ 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
@@ -74,6 +75,7 @@ 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
@@ -129,6 +131,7 @@ 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
@@ -177,6 +180,7 @@ 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
@@ -232,6 +236,7 @@ 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
@@ -280,6 +285,7 @@ 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 ba60000..32281c9 100644 (file)
--- 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 7f040000 00000000
+// OBJDUMP-NEXT: 0070 015001e4 1f0f007f 7f0c0000 00000000
 // special_sgpr
 // OBJDUMP-NEXT: 0080 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 0090 00000000 00000000 00000000 00000000
@@ -91,6 +91,7 @@ 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
@@ -134,6 +135,7 @@ 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 7f885b4..ca0783f 100644 (file)
--- 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 5e040000 00000000
+// OBJDUMP-NEXT: 0070 015001e4 130f007f 5e0c0000 00000000
 // special_sgpr
 // OBJDUMP-NEXT: 0080 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 0090 00000000 00000000 00000000 00000000
@@ -87,6 +87,7 @@ 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
@@ -126,6 +127,7 @@ 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 6ab70e5..4c10763 100644 (file)
--- 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 7f000000 00000000
+// OBJDUMP-NEXT: 0070 c1500104 1f0f007f 7f080000 00000000
 
 .text
 // ASM: .text
@@ -77,6 +77,7 @@ 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
@@ -117,6 +118,7 @@ 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 bf2e608..46a7303 100644 (file)
--- 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 5e000000 00000000
+// OBJDUMP-NEXT: 0070 01510104 130f007f 5e080000 00000000
 
 .text
 // ASM: .text
@@ -75,6 +75,7 @@ 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
@@ -112,6 +113,7 @@ 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 7da24a9..771f37d 100644 (file)
--- a/llvm/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s
+++ b/llvm/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s
@@ -39,6 +39,7 @@
 ; 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:   
@@ -65,6 +66,7 @@
       .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 602b456..6cdd215 100644 (file)
--- a/llvm/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s
+++ b/llvm/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s
@@ -20,6 +20,7 @@
 // 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
@@ -51,6 +52,7 @@
       .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 cad60ea..064ccb3 100644 (file)
--- 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,6 +13,7 @@
 // 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
@@ -32,6 +33,7 @@
       .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 9f85498..d9dbd59 100644 (file)
--- 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 7f000000 00000000
+// OBJDUMP-NEXT: 0070 c2500104 1f0f007f 7f080000 00000000
 // special_sgpr
 // OBJDUMP-NEXT: 0080 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 0090 00000000 00000000 00000000 00000000
@@ -101,6 +101,7 @@ 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
@@ -140,6 +141,7 @@ 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
@@ -259,6 +261,7 @@ 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
@@ -269,6 +272,7 @@ 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
@@ -286,6 +290,7 @@ 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
@@ -296,6 +301,7 @@ 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:

diff --git a/llvm/test/MC/AMDGPU/hsa-v4.s b/llvm/test/MC/AMDGPU/hsa-v4.s
index 6a824b8..79dd9a0 100644 (file)
--- 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 7f000000 00000000
+// OBJDUMP-NEXT: 0070 c2500104 1f0f007f 7f080000 00000000
 // special_sgpr
 // OBJDUMP-NEXT: 0080 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 0090 00000000 00000000 00000000 00000000
@@ -102,6 +102,7 @@ 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,6 +142,7 @@ 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
@@ -260,6 +262,7 @@ 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
@@ -269,6 +272,7 @@ 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
@@ -287,6 +291,7 @@ 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
@@ -297,6 +302,7 @@ 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:

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 73022c3..a84b2ab 100644 (file)
--- 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: 0x110
+#LLVM-NEXT:      Size: 0x128
 #LLVM-NEXT:      Note {
 #LLVM-NEXT:        Owner: AMDGPU
-#LLVM-NEXT:        Data size: 0xFC
+#LLVM-NEXT:        Data size: 0x111
 #LLVM-NEXT:        Type: NT_AMDGPU_METADATA (AMDGPU Metadata)
 #LLVM-NEXT:        AMDGPU Metadata: ---
 #LLVM-NEXT:  amdhsa.kernels:
@@ -25,6 +25,7 @@
 #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:
@@ -38,7 +39,7 @@
 
 # GNU:      Displaying notes found in: .note
 # GNU-NEXT:   Owner                Data size        Description
-# GNU-NEXT:   AMDGPU               0x000000fc       NT_AMDGPU_METADATA (AMDGPU Metadata)
+# GNU-NEXT:   AMDGPU               0x00000111       NT_AMDGPU_METADATA (AMDGPU Metadata)
 # GNU-NEXT:     AMDGPU Metadata:
 # GNU-NEXT:         ---
 # GNU-NEXT: amdhsa.kernels:
@@ -50,6 +51,7 @@
 # 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:
@@ -69,6 +71,7 @@
       .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

diff --git a/llvm/test/tools/llvm-readobj/ELF/note-amdgpu.test b/llvm/test/tools/llvm-readobj/ELF/note-amdgpu.test
index f2d89d1..5d4f41c 100644 (file)
--- a/llvm/test/tools/llvm-readobj/ELF/note-amdgpu.test
+++ b/llvm/test/tools/llvm-readobj/ELF/note-amdgpu.test
@@ -4,7 +4,7 @@
 
 # GNU:      Displaying notes found in: .note.foo
 # GNU-NEXT:   Owner                Data size        Description
-# GNU-NEXT:   AMDGPU               0x000000e6       NT_AMDGPU_METADATA (AMDGPU Metadata)
+# GNU-NEXT:   AMDGPU               0x000000fb       NT_AMDGPU_METADATA (AMDGPU Metadata)
 # GNU-NEXT:     AMDGPU Metadata:
 # GNU-NEXT:         ---
 # GNU-NEXT: amdhsa.kernels:
@@ -16,6 +16,7 @@
 # GNU-NEXT:     .private_segment_fixed_size: 3
 # GNU-NEXT:     .sgpr_count:     6
 # GNU-NEXT:     .symbol:         foo
+# GNU-NEXT:     .uses_dynamic_stack: true
 # GNU-NEXT:     .vgpr_count:     7
 # GNU-NEXT:     .wavefront_size: 5
 # GNU-NEXT: amdhsa.version:
@@ -37,7 +38,7 @@
 # LLVM-NEXT:     Size:
 # LLVM-NEXT:     Note {
 # LLVM-NEXT:       Owner: AMDGPU
-# LLVM-NEXT:       Data size: 0xE6
+# LLVM-NEXT:       Data size: 0xFB
 # LLVM-NEXT:       Type: NT_AMDGPU_METADATA (AMDGPU Metadata)
 # LLVM-NEXT:       AMDGPU Metadata: ---
 # LLVM-NEXT: amdhsa.kernels:
@@ -49,6 +50,7 @@
 # LLVM-NEXT:     .private_segment_fixed_size: 3
 # LLVM-NEXT:     .sgpr_count:     6
 # LLVM-NEXT:     .symbol:         foo
+# LLVM-NEXT:     .uses_dynamic_stack: true
 # LLVM-NEXT:     .vgpr_count:     7
 # LLVM-NEXT:     .wavefront_size: 5
 # LLVM-NEXT: amdhsa.version:
@@ -60,7 +62,7 @@
 # LLVM-NEXT:   }
 # LLVM-NEXT:   NoteSection {
 # LLVM-NEXT:     Name: .note.unknown
-# LLVM-NEXT:     Offset: 0x13C
+# LLVM-NEXT:     Offset: 0x150
 # LLVM-NEXT:     Size: 0x18
 # LLVM-NEXT:     Note {
 # LLVM-NEXT:       Owner: AMDGPU
@@ -87,12 +89,17 @@
 #       .kernarg_segment_size: 1
 #       .group_segment_fixed_size: 2
 #       .private_segment_fixed_size: 3
+#       .uses_dynamic_stack: true
 #       .kernarg_segment_align: 4
 #       .wavefront_size: 5
 #       .sgpr_count: 6
 #       .vgpr_count: 7
 #       .max_flat_workgroup_size: 8
 # .end_amdgpu_metadata
+#
+## Here's one way to get the contents of .note.foo in the test input from %t.o:
+# $ llvm-objcopy -O binary --only-section=.note %t.o note.out
+# $ xxd -p note.out | tr -d '\n' | tr a-z A-Z
 
 --- !ELF
 FileHeader:
@@ -102,7 +109,7 @@ FileHeader:
 Sections:
   - Name:        .note.foo
     Type:        SHT_NOTE
-    Content:     07000000E600000020000000414D44475055000082AE616D646873612E6B65726E656C73918AB92E67726F75705F7365676D656E745F66697865645F73697A6502B62E6B65726E6172675F7365676D656E745F616C69676E04B52E6B65726E6172675F7365676D656E745F73697A6501B82E6D61785F666C61745F776F726B67726F75705F73697A6508A52E6E616D65A3666F6FBB2E707269766174655F7365676D656E745F66697865645F73697A6503AB2E736770725F636F756E7406A72E73796D626F6CA3666F6FAB2E766770725F636F756E7407AF2E7761766566726F6E745F73697A6505AE616D646873612E76657273696F6E9201000000
+    Content:     07000000FB00000020000000414D44475055000082AE616D646873612E6B65726E656C73918BB92E67726F75705F7365676D656E745F66697865645F73697A6502B62E6B65726E6172675F7365676D656E745F616C69676E04B52E6B65726E6172675F7365676D656E745F73697A6501B82E6D61785F666C61745F776F726B67726F75705F73697A6508A52E6E616D65A3666F6FBB2E707269766174655F7365676D656E745F66697865645F73697A6503AB2E736770725F636F756E7406A72E73796D626F6CA3666F6FB32E757365735F64796E616D69635F737461636BC3AB2E766770725F636F756E7407AF2E7761766566726F6E745F73697A6505AE616D646873612E76657273696F6E92010000
   - Name:        .note.unknown
     Type:        SHT_NOTE
     Notes: