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

[abinavpp]

Summary

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.

[arsenm]

What's broken in the runtime? If it didn't read it before, how can this break something?

llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

420–422

Swap order of checks

[abinavpp]

Address review comment

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

877–879

This isn't querying the module for the code object version; probably should have that in the asm printer doInitialization

[abinavpp]

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

877–879

I didn't get you. AMDGPUAsmPrinter's ctor checks for the code-obj-ver using isHsaAbiVersion*, which checks for the AmdhsaCodeObjectVersion like this. Or should we override emitKernel() for MetadataStreamerMsgPackV5 and set this entry there. I don't think I understood the problem that we're trying to fixing by overriding AsmPrinter's doInitialization().

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

877–879

@cfang had a patch to switch to querying the code object version from the IR module instead of the global command line flag

[abinavpp]

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

877–879

I think we should do this in a separate patch and for now, stick to getAmdhsaCodeObjectVersion() in this patch for consistency. I believe we should replace all getAmdhsaCodeObjectVersion() queries with the "amdgpu_code_object_version" module flag query. I would like to see that metadata map to something like the enum CodeObjectVersionKind in clang's include/clang/Basic/TargetOptions.h in llvm. This will be a big change I guess. I couldn't find @cfang 's patch for this, I'm assuming he's working on it or is done with it.

[abinavpp]

Oops, I think I broke docs-llvm-html: https://lab.llvm.org/buildbot/#/builders/30/builds/27004. Sorry, I don't think I enabled this in my cmake at the moment. Will this fix the error:

diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index da8a5cf19dbc..eac0883c4949 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -3577,6 +3577,7 @@ Code object V5 metadata is the same as
      ".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

[abinavpp]

Checked by running sphinx manually. Fix: bb4abb1a005c11add7f3206faf0f526b2e54e6f8

llvm/docs/AMDGPUUsage.rst

".private_segment_fixed_size"       integer        Required  The amount of fixed
                                                                   space must be added
                                                                   to this value for the
                                                                   call stack.
      ".kernarg_segment_align"            integer        Required  The maximum byte
                                                                   alignment of
                                                                   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
                                                                   registers required by a
                                                                   wavefront for
                                                                   GFX6-GFX9. A register
                                                                   is required if it is
                                                                   used explicitly, or

 
 .. warning::
   Code object V5 is not the default code object version emitted by this version
   of LLVM.
 
 
 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
      :name: amdgpu-amdhsa-code-object-metadata-map-table-v5
 
      ================= ============== ========= =======================================
      String Key        Value Type     Required? Description
      ================= ============== ========= =======================================
      "amdhsa.version"  sequence of    Required  - The first integer is the major
                        2 integers                 version. Currently 1.
                                                 - The second integer is the minor
                                                   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
      :name: amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v5
 
      ====================== ============== ========= ================================
      String Key             Value Type     Required? Description
      ====================== ============== ========= ================================
      ".value_kind"          string         Required  Kernel argument kind that
                                                      specifies how to set up the
458     1 bit   ENABLE_WAVEFRONT_SIZE32         GFX6-GFX9
                                                        - If 0 execute in
                                                          wavefront size 64 mode.
                                                        - If 1 execute in
                                                          native wavefront size
                                                          32 mode.
      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.
      468     1 bit   RESERVED_468                    Deprecated, must be 0.
      469:471 3 bits                                  Reserved, must be 0.
      511:472 5 bytes                                 Reserved, must be 0.
      512     **Total size 64 bytes.**
      ======= ====================================================================

llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

if (MFI.hasFlatScratchInit()) {
     KernelCodeProperties |=
         amdhsa::KERNEL_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT;
   }
   if (MF.getSubtarget<GCNSubtarget>().isWave32()) {
     KernelCodeProperties |=
         amdhsa::KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32;
   }
 
-  if (CurrentProgramInfo.DynamicCallStack) {
+  if (CurrentProgramInfo.DynamicCallStack &&
+      AMDGPU::getAmdhsaCodeObjectVersion() >= 5) {
     KernelCodeProperties |= amdhsa::KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK;
   }
 
   return KernelCodeProperties;
 }
 
 amdhsa::kernel_descriptor_t AMDGPUAsmPrinter::getAmdhsaKernelDescriptor(
     const MachineFunction &MF,
     const SIProgramInfo &PI) const {

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

msgpack::MapDocNode MetadataStreamerMsgPackV3::getHSAKernelProps(
 
   Align MaxKernArgAlign;
   Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
       STM.getKernArgSegmentSize(F, MaxKernArgAlign));
   Kern[".group_segment_fixed_size"] =
       Kern.getDocument()->getNode(ProgramInfo.LDSSize);
   Kern[".private_segment_fixed_size"] =
       Kern.getDocument()->getNode(ProgramInfo.ScratchSize);
+  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"] =
       Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value());
   Kern[".wavefront_size"] =
       Kern.getDocument()->getNode(STM.getWavefrontSize());
   Kern[".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumSGPR);
   Kern[".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumVGPR);

llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp

case amdhsa::KERNEL_CODE_PROPERTIES_OFFSET:
     if (isGFX9() &&
         (TwoByteBuffer & KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32)) {
       return MCDisassembler::Fail;
     } else if (isGFX10Plus()) {
       PRINT_DIRECTIVE(".amdhsa_wavefront_size32",
                       KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32);
     }
 
+    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;
 
     return MCDisassembler::Success;
 
   case amdhsa::RESERVED2_OFFSET:
     // 6 bytes from here are reserved, must be 0.

llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp

PRINT_FIELD(OS, ".amdhsa_user_sgpr_flat_scratch_init", KD,
                 amdhsa::KERNEL_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT);
   PRINT_FIELD(OS, ".amdhsa_user_sgpr_private_segment_size", KD,
               kernel_code_properties,
               amdhsa::KERNEL_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE);
   if (IVersion.Major >= 10)
     PRINT_FIELD(OS, ".amdhsa_wavefront_size32", KD,
                 kernel_code_properties,
                 amdhsa::KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32);
+  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"
                    : ".amdhsa_system_sgpr_private_segment_wavefront_offset"),
               KD, compute_pgm_rsrc2,
               amdhsa::COMPUTE_PGM_RSRC2_ENABLE_PRIVATE_SEGMENT);
   PRINT_FIELD(OS, ".amdhsa_system_sgpr_workgroup_id_x", KD,
               compute_pgm_rsrc2,

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

 ; WORKAROUND: .amdhsa_user_sgpr_count 15
 ; NOWORKAROUND: .amdhsa_user_sgpr_count 0
 ; GCN-NEXT: .amdhsa_user_sgpr_dispatch_ptr 0
 ; GCN-NEXT: .amdhsa_user_sgpr_queue_ptr 0
 ; GCN-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 0
 ; 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
 ; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_z 0
 ; GCN-NEXT: .amdhsa_system_sgpr_workgroup_info 0
 ; GCN-NEXT: .amdhsa_system_vgpr_workitem_id 0
 ; WORKAROUND: ; COMPUTE_PGM_RSRC2:USER_SGPR: 15
 ; NOWORKAROUND: ; COMPUTE_PGM_RSRC2:USER_SGPR: 0

 ; WORKAROUND: .amdhsa_user_sgpr_count 15
 ; NOWORKAROUND: .amdhsa_user_sgpr_count 0
 ; GCN-NEXT: .amdhsa_user_sgpr_dispatch_ptr 0
 ; GCN-NEXT: .amdhsa_user_sgpr_queue_ptr 0
 ; GCN-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 0
 ; 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
 ; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_z 0
 ; GCN-NEXT: .amdhsa_system_sgpr_workgroup_info 0
 ; GCN-NEXT: .amdhsa_system_vgpr_workitem_id 0
 ; WORKAROUND: ; COMPUTE_PGM_RSRC2:USER_SGPR: 15
 ; NOWORKAROUND: ; COMPUTE_PGM_RSRC2:USER_SGPR: 0

 ; WORKAROUND: .amdhsa_user_sgpr_count 15
 ; NOWORKAROUND: .amdhsa_user_sgpr_count 2
 ; GCN-NEXT: .amdhsa_user_sgpr_dispatch_ptr 0
 ; GCN-NEXT: .amdhsa_user_sgpr_queue_ptr 1
 ; GCN-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 0
 ; 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
 ; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_z 0
 ; GCN-NEXT: .amdhsa_system_sgpr_workgroup_info 0
 ; GCN-NEXT: .amdhsa_system_vgpr_workitem_id 0
 ; WORKAROUND: ; COMPUTE_PGM_RSRC2:USER_SGPR: 15
 ; NOWORKAROUND: ; COMPUTE_PGM_RSRC2:USER_SGPR: 2

 ; WORKAROUND: .amdhsa_user_sgpr_count 13
 ; NOWORKAROUND: .amdhsa_user_sgpr_count 8
 ; GCN-NEXT: .amdhsa_user_sgpr_dispatch_ptr 1
 ; GCN-NEXT: .amdhsa_user_sgpr_queue_ptr 1
 ; GCN-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 1
 ; 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
 ; GCN-NEXT: .amdhsa_system_sgpr_workgroup_id_z 1
 ; GCN-NEXT: .amdhsa_system_sgpr_workgroup_info 0
 ; GCN-NEXT: .amdhsa_system_vgpr_workitem_id 0
 ; WORKAROUND: ; COMPUTE_PGM_RSRC2:USER_SGPR: 13
 ; NOWORKAROUND: ; COMPUTE_PGM_RSRC2:USER_SGPR: 8

llvm/test/CodeGen/AMDGPU/recursion.ll

 
 ; For an arbitrary recursive call, report a large number for unknown stack
 ; usage for code object v4 and older
 ; CHECK-LABEL: {{^}}calls_recursive:
 ; CHECK: .amdhsa_private_segment_fixed_size 16400{{$}}
 ;
 ; 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
 }
 
 ; Make sure we do not report a huge stack size for tail recursive
 ; functions
 ; CHECK-LABEL: {{^}}kernel_indirectly_calls_tail_recursive:
 ; CHECK: .amdhsa_private_segment_fixed_size 0{{$}}
 define amdgpu_kernel void @kernel_indirectly_calls_tail_recursive() {
   call void @calls_tail_recursive()
   ret void
 }
 
 ; TODO: Even though tail_recursive is only called as a tail call, we
 ; end up treating it as generally recursive call from the regular call
 ; in the kernel.
 
 ; CHECK-LABEL: {{^}}kernel_calls_tail_recursive:
 ; CHECK: .amdhsa_private_segment_fixed_size 16384{{$}}
 ;
 ; 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
 }
 
 ; CHECK-LABEL: {{^}}kernel_calls_tail_recursive_with_stack:
 ; CHECK: .amdhsa_private_segment_fixed_size 16384{{$}}
 ;
 ; 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
 }

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

 ; 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.
 
 @gv.fptr0 = external hidden unnamed_addr addrspace(4) constant void()*, align 4
 
 ; GCN-LABEL: unreachable:
 ; Function info:
 ; codeLenInByte = 4
 define internal fastcc void @unreachable() {
   %fptr = load void()*, void()* addrspace(4)* @gv.fptr0
   call void %fptr()
   unreachable
 }
 
 
 ; GCN-LABEL: entry:
 ; GCN-NOT: s_swappc_b64
 ; 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
 
 bb1:
   tail call fastcc void @unreachable()
   unreachable
 
 bb2:
   ret void
 }

llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll

 ; VI-NEXT:     .amdhsa_user_sgpr_count 6
 ; VI-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; VI-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 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
 ; VI-NEXT:     .amdhsa_system_sgpr_workgroup_id_z 0
 ; VI-NEXT:     .amdhsa_system_sgpr_workgroup_info 0
 ; VI-NEXT:     .amdhsa_system_vgpr_workitem_id 0
 ; VI-NEXT:     .amdhsa_next_free_vgpr 1
 ; VI-NEXT:     .amdhsa_next_free_sgpr 8

 ; GFX9-NEXT:     .amdhsa_user_sgpr_count 6
 ; GFX9-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; GFX9-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 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
 ; GFX9-NEXT:     .amdhsa_system_sgpr_workgroup_id_z 0
 ; GFX9-NEXT:     .amdhsa_system_sgpr_workgroup_info 0
 ; GFX9-NEXT:     .amdhsa_system_vgpr_workitem_id 0
 ; GFX9-NEXT:     .amdhsa_next_free_vgpr 1
 ; GFX9-NEXT:     .amdhsa_next_free_sgpr 8

 ; VI-NEXT:     .amdhsa_user_sgpr_count 6
 ; VI-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; VI-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 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
 ; VI-NEXT:     .amdhsa_system_sgpr_workgroup_id_z 0
 ; VI-NEXT:     .amdhsa_system_sgpr_workgroup_info 0
 ; VI-NEXT:     .amdhsa_system_vgpr_workitem_id 0
 ; VI-NEXT:     .amdhsa_next_free_vgpr 1
 ; VI-NEXT:     .amdhsa_next_free_sgpr 8

 ; GFX9-NEXT:     .amdhsa_user_sgpr_count 6
 ; GFX9-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; GFX9-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 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
 ; GFX9-NEXT:     .amdhsa_system_sgpr_workgroup_id_z 0
 ; GFX9-NEXT:     .amdhsa_system_sgpr_workgroup_info 0
 ; GFX9-NEXT:     .amdhsa_system_vgpr_workitem_id 0
 ; GFX9-NEXT:     .amdhsa_next_free_vgpr 1
 ; GFX9-NEXT:     .amdhsa_next_free_sgpr 8

 ; VI-NEXT:     .amdhsa_user_sgpr_count 6
 ; VI-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; VI-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 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
 ; VI-NEXT:     .amdhsa_system_sgpr_workgroup_id_z 0
 ; VI-NEXT:     .amdhsa_system_sgpr_workgroup_info 0
 ; VI-NEXT:     .amdhsa_system_vgpr_workitem_id 0
 ; VI-NEXT:     .amdhsa_next_free_vgpr 1
 ; VI-NEXT:     .amdhsa_next_free_sgpr 8

 ; GFX9-NEXT:     .amdhsa_user_sgpr_count 6
 ; GFX9-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; GFX9-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 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
 ; GFX9-NEXT:     .amdhsa_system_sgpr_workgroup_id_z 0
 ; GFX9-NEXT:     .amdhsa_system_sgpr_workgroup_info 0
 ; GFX9-NEXT:     .amdhsa_system_vgpr_workitem_id 0
 ; GFX9-NEXT:     .amdhsa_next_free_vgpr 1
 ; GFX9-NEXT:     .amdhsa_next_free_sgpr 8

llvm/test/MC/AMDGPU/hsa-gfx10-v3.s

 // OBJDUMP-NEXT: 0000 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 0010 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 0020 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 0030 0000ac60 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 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
 // OBJDUMP-NEXT: 00a0 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 00b0 00000060 80000000 00000000 00000000
 
 .text
 // ASM: .text
.amdhsa_kernel complete
   .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_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
   .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

 // 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_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
 // 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

llvm/test/MC/AMDGPU/hsa-gfx11-v3.s

 // OBJDUMP-NEXT: 0000 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 0010 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 0020 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 0030 0000ac60 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 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
 // OBJDUMP-NEXT: 00a0 00000000 00000000 00000000 00000000
 // OBJDUMP-NEXT: 00b0 00000060 80000000 00000000 00000000
 
 .text
 // ASM: .text
.amdhsa_kernel complete
   .amdhsa_private_segment_fixed_size 1
   .amdhsa_kernarg_size 8
   .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_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
   .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

 // ASM-NEXT: .amdhsa_kernarg_size 8
 // ASM-NEXT: .amdhsa_user_sgpr_count 9
 // 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_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
 // 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

llvm/test/MC/AMDGPU/hsa-gfx90a-v3.s

 // 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 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
 
 .amdgcn_target "amdgcn-amd-amdhsa--gfx90a+xnack+sram-ecc"
 // ASM: .amdgcn_target "amdgcn-amd-amdhsa--gfx90a+xnack+sram-ecc"
 
 .p2align 8
.amdhsa_kernel complete
   .amdhsa_private_segment_fixed_size 1
   .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

 // 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

llvm/test/MC/AMDGPU/hsa-gfx940-v3.s

 // 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 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
 
 .amdgcn_target "amdgcn-amd-amdhsa--gfx940+xnack+sram-ecc"
 // ASM: .amdgcn_target "amdgcn-amd-amdhsa--gfx940+xnack+sram-ecc"
 
 .p2align 8

 .amdhsa_kernel complete
   .amdhsa_group_segment_fixed_size 1
   .amdhsa_private_segment_fixed_size 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_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
   .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

 // ASM-NEXT: .amdhsa_private_segment_fixed_size 1
 // ASM-NEXT: .amdhsa_kernarg_size 0
 // ASM-NEXT: .amdhsa_user_sgpr_count 9
 // 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_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
 // 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

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

 ; CHECK-NEXT:     .language_version: 
 ; CHECK-NEXT:       - 2
 ; CHECK-NEXT:       - 0
 ; CHECK-NEXT:     .max_flat_workgroup_size: 256
 ; CHECK-NEXT:     .name:           test_kernel
 ; 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:   
 ; CHECK-NEXT:   - '1:1:4:%d\n'
 ; CHECK-NEXT:   - '2:1:8:%g\n'
 ; CHECK-NEXT: amdhsa.version:  
 ; CHECK-NEXT:   - 1
 ; CHECK-NEXT:   - 0
- .name:            test_kernel
       .symbol:      test_kernel@kd
       .language:        OpenCL C
       .language_version:
         - 2
         - 0
       .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
       .args:
         - .type_name:      char
           .size:          1

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

 // CHECK:          .name:           test_kernel
 // CHECK:          .private_segment_fixed_size: 32
 // CHECK:          .reqd_workgroup_size: 
 // CHECK-NEXT:       - 1
 // CHECK-NEXT:       - 2
 // 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
 // CHECK:          .workgroup_size_hint: 
 // CHECK-NEXT:       - 8
 // CHECK-NEXT:       - 16
 // CHECK-NEXT:       - 32
 // CHECK:      amdhsa.printf:   
- .name:            test_kernel
       .symbol:      test_kernel@kd
       .language:        OpenCL C
       .language_version:
         - 2
         - 0
       .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
       .reqd_workgroup_size:
         - 1
         - 2
         - 4
       .workgroup_size_hint:
         - 8
         - 16
         - 32
       .vec_type_hint:       int
 .end_amdgpu_metadata

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

 // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK %s
 // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK %s
 // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK %s
 
 // CHECK:       .amdgpu_metadata
 // CHECK:      amdhsa.kernels:  
 // CHECK:        - .group_segment_fixed_size: 24
 // CHECK:          .kernarg_segment_align: 16
 // CHECK:          .kernarg_segment_size: 24
 // CHECK:          .max_flat_workgroup_size: 256
 // CHECK:          .name:           test_kernel
 // CHECK:          .private_segment_fixed_size: 16
 // 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
 // CHECK:      amdhsa.version:  
 // CHECK-NEXT:   - 1
 // CHECK-NEXT:   - 0
 .amdgpu_metadata
   amdhsa.version:
     - 1
     - 0
   amdhsa.printf:
     - '1:1:4:%d\n'
     - '2:1:8:%g\n'
   amdhsa.kernels:
     - .name:            test_kernel
       .symbol:      test_kernel@kd
       .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
       .sgpr_count:               40
       .vgpr_count:               14
       .sgpr_spill_count:         1
       .vgpr_spill_count:         1
 .end_amdgpu_metadata

llvm/test/MC/AMDGPU/hsa-v3.s

 // 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
+// OBJDUMP-NEXT: 0070 c2500104 1f0f007f 7f000000 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
.amdhsa_kernel complete
   .amdhsa_kernarg_size 8
   .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

 // 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
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

llvm/test/MC/AMDGPU/hsa-v4.s

 // 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
+// OBJDUMP-NEXT: 0070 c2500104 1f0f007f 7f000000 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
.amdhsa_kernel complete
   .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

 // 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
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
-      .uses_dynamic_stack: true
       .private_segment_fixed_size: 32
       .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

llvm/test/MC/AMDGPU/hsa-v5-uses-dynamic-stack.s

-// RUN: llvm-mc -triple amdgcn-amd-amdhsa -mcpu=gfx904 --amdhsa-code-object-version=4 -mattr=+xnack < %s | FileCheck --check-prefix=ASM %s
-// RUN: llvm-mc -triple amdgcn-amd-amdhsa -mcpu=gfx904 --amdhsa-code-object-version=4 -mattr=+xnack -filetype=obj < %s > %t
+// 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
- .name:       amd_kernel_code_t_test_all
       .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
-      .uses_dynamic_stack: true
       .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

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

 ## This test is checking the handling of valid note entries for AMDGPU code
 ## object v3.
 
 # REQUIRES: amdgpu-registered-target
 
 # RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj < %s | llvm-readobj --notes - | FileCheck %s --match-full-lines --check-prefix=LLVM
 # RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj < %s | llvm-readelf --notes - | FileCheck %s --match-full-lines --check-prefix=GNU
 
 #LLVM:       Notes [
 #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:
 #LLVM-NEXT:    - .group_segment_fixed_size: 16
 #LLVM-NEXT:      .kernarg_segment_align: 64
 #LLVM-NEXT:      .kernarg_segment_size: 8
 #LLVM-NEXT:      .max_flat_workgroup_size: 256
 #LLVM-NEXT:      .name:           test_kernel
 #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:
 #LLVM-NEXT:    - 1
 #LLVM-NEXT:    - 0
 #LLVM-NEXT:  ...
 #LLVM-EMPTY:
 #LLVM-NEXT:      }
 #LLVM-NEXT:    }
 #LLVM-NEXT:  ]
 
 # 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:
 # GNU-NEXT:   - .group_segment_fixed_size: 16
 # GNU-NEXT:     .kernarg_segment_align: 64
 # GNU-NEXT:     .kernarg_segment_size: 8
 # GNU-NEXT:     .max_flat_workgroup_size: 256
 # GNU-NEXT:     .name:           test_kernel
 # 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:
 # GNU-NEXT:   - 1
 # GNU-NEXT:   - 0
 # GNU-NEXT: ...
 
 .amdgpu_metadata
   amdhsa.version:
     - 1
     - 0
   amdhsa.kernels:
     - .name:                       test_kernel
       .symbol:                     test_kernel@kd
       .group_segment_fixed_size:   16
       .kernarg_segment_align:      64
       .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
 .end_amdgpu_metadata