[AMDGPU] Add the uses_dynamic_stack field to the kernel descriptor and the kernel metadata map

[abinavpp]

Summary

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.

[kzhuravl]

Do we want to tie this bit to a particular code object version? Code object v4 and up? Why did you pick code object v3 and up?

llvm/docs/AMDGPUUsage.rst

4006

If it is "Reserved, must be 0.", then compiler does not set it. But you clearly set it in in AsmPrinter

llvm/include/llvm/Support/AMDHSAKernelDescriptor.h

165

Why not eat up the reserved bit at index 11 and decrease number of reserved bits to 4?

[arsenm]

llvm/include/llvm/Support/AMDHSAKernelDescriptor.h

181–182

I thought this was already defined here, which was the reason to set it. I'm not sure it makes sense to add the field if it's not already there

[arsenm]

llvm/docs/AMDGPUUsage.rst

4003–4006

Since we are actually re-introducing this and not recycling an existing field, we should take the chance to rename it. The call part of the name is misleading since we can need this in situations that do not have calls. How about just DYNAMIC_STACK_SIZE?

llvm/include/llvm/Support/AMDHSAKernelDescriptor.h

181–182

We apparently do need to place this here since rocr doesn't want to inspect metadata

[abinavpp]

In D128344#3601894, @kzhuravl wrote:

Do we want to tie this bit to a particular code object version? Code object v4 and up? Why did you pick code object v3 and up?

My intention is to add this to the kernel descriptor after seeing that rocr relies on it. My understanding is that modifying kernel descriptor affects code object v3 and above. Should we only set this for v4 and above? If so, how should we modify amdgpu-amdhsa-kernel-descriptor-v3-table in AMDGPUUsage.rst? I'm not sure about adding v4+ specific fields there.

llvm/docs/AMDGPUUsage.rst

4003–4006

If there's no requirement to keep the naming consistent with amd_kernel_code_t, we should rename this. DYNAMIC_STACK_SIZE sounds like an integer field, how about IS_DYNAMIC_STACK?

4006

Sorry, IS_DYNAMIC_CALLSTACK should be at the 468th bit since we have KERNEL_CODE_PROPERTY(IS_DYNAMIC_CALLSTACK, 20, 1). I'll fix this once we agree on the bit position for this field.

llvm/include/llvm/Support/AMDHSAKernelDescriptor.h

165

I did that initially, but changed it to the 20th bit to keep things consistent after noticing the "Must be kept backwards compatible." comment and that these fields were at the same bit position as the ones in amd_kernel_code_t's code properties field (i.e. amd_code_property_mask_t of AMDKernelCodeT.h). My understanding is limited here, but do we need to make sure we use the same bit position as that of v2's amd_kernel_code_t? If not, then we should just just use the first bit of RESEREVED1 and avoid the extension of code property to 32 bits. Also, was there any relevant reason for keeping the bit positions conform to that of amd_kernel_code_t to date?

[abinavpp]

• As per internal discussion, use the 11th bit
• Rename the field to is_dynamic_stack

[b-sumner]

The "is_" prefix is not helpful. I suggest either "uses_dynamic_stack" or just "dynamic_stack".

[abinavpp]

Rename the field to uses_dynamic_stack

[arsenm]

I see the kernel descriptor changes reflected in the test updates, but I don't see the metadata entries tested

[abinavpp]

In D128344#3615280, @arsenm wrote:

I see the kernel descriptor changes reflected in the test updates, but I don't see the metadata entries tested

By "metadata entries", do you mean an entry in the kernel metadata map under NT_AMDGPU_METADATA? I didn't add the corresponding field there.

[arsenm]

In D128344#3616164, @abinavpp wrote:

In D128344#3615280, @arsenm wrote:

I see the kernel descriptor changes reflected in the test updates, but I don't see the metadata entries tested

By "metadata entries", do you mean an entry in the kernel metadata map under NT_AMDGPU_METADATA? I didn't add the corresponding field there.

Yes. It has .private_segment_fixed_size, so should also have this for consistency

[abinavpp]

• Rebase
• Add the field to the kernel metadata map

[arsenm]

LGTM

[abinavpp]

• Rebase
• Make the field optional

[arsenm]

Could use some assembler tests demonstrating the default behavior

[abinavpp]

In D128344#3630657, @arsenm wrote:

Could use some assembler tests demonstrating the default behavior

The objdump assertions of the "minimal" symbol in hsa-v3.s and hsa-v4.s under MC/AMDGPU checks the default bits in the kernel descriptor.

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
.. table:: Code Object V3 Kernel Descriptor
      458     1 bit   ENABLE_WAVEFRONT_SIZE32         GFX6-GFX9
                                                        Reserved, must be 0.
                                                      GFX10-GFX11
                                                        - If 0 execute in
                                                          wavefront size 64 mode.
                                                        - 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.
      469:471 3 bits                                  Reserved, must be 0.
      511:472 5 bytes                                 Reserved, must be 0.
      512     **Total size 64 bytes.**
      ======= ====================================================================
 
``.amdhsa_user_sgpr_flat_scratch_init``                  0                   GFX6-GFX10   Controls ENABLE_SGPR_FLAT_SCRATCH_INIT in
                                                                                   (except      :ref:`amdgpu-amdhsa-kernel-descriptor-v3-table`.
                                                                                   GFX940)
      ``.amdhsa_user_sgpr_private_segment_size``               0                   GFX6-GFX11   Controls ENABLE_SGPR_PRIVATE_SEGMENT_SIZE in
                                                                                                :ref:`amdgpu-amdhsa-kernel-descriptor-v3-table`.
      ``.amdhsa_wavefront_size32``                             Target              GFX10-GFX11  Controls ENABLE_WAVEFRONT_SIZE32 in
                                                               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)
      ``.amdhsa_enable_private_segment``                       0                   GFX940,      Controls ENABLE_PRIVATE_SEGMENT in
                                                                                   GFX11        :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx11-table`.
      ``.amdhsa_system_sgpr_workgroup_id_x``                   1                   GFX6-GFX11   Controls ENABLE_SGPR_WORKGROUP_ID_X in
                                                                                                :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx11-table`.
      ``.amdhsa_system_sgpr_workgroup_id_y``                   0                   GFX6-GFX11   Controls ENABLE_SGPR_WORKGROUP_ID_Y in

llvm/include/llvm/Support/AMDHSAKernelDescriptor.h

enum : int32_t {
   KERNEL_CODE_PROPERTY(ENABLE_SGPR_DISPATCH_PTR, 1, 1),
   KERNEL_CODE_PROPERTY(ENABLE_SGPR_QUEUE_PTR, 2, 1),
   KERNEL_CODE_PROPERTY(ENABLE_SGPR_KERNARG_SEGMENT_PTR, 3, 1),
   KERNEL_CODE_PROPERTY(ENABLE_SGPR_DISPATCH_ID, 4, 1),
   KERNEL_CODE_PROPERTY(ENABLE_SGPR_FLAT_SCRATCH_INIT, 5, 1),
   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
 
 // Kernel descriptor. Must be kept backwards compatible.
 struct kernel_descriptor_t {
   uint32_t group_segment_fixed_size;
   uint32_t private_segment_fixed_size;
   uint32_t kernarg_size;
   uint8_t reserved0[4];
   int64_t kernel_code_entry_byte_offset;
   uint8_t reserved1[20];
   uint32_t compute_pgm_rsrc3; // GFX10+ and GFX90A+
   uint32_t compute_pgm_rsrc1;
   uint32_t compute_pgm_rsrc2;
   uint16_t kernel_code_properties;
   uint8_t reserved2[6];
 };
 
 enum : uint32_t {
   GROUP_SEGMENT_FIXED_SIZE_OFFSET = 0,
   PRIVATE_SEGMENT_FIXED_SIZE_OFFSET = 4,
   KERNARG_SIZE_OFFSET = 8,
   RESERVED0_OFFSET = 12,
   KERNEL_CODE_ENTRY_BYTE_OFFSET_OFFSET = 16,
   RESERVED1_OFFSET = 24,

llvm/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp

if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false,
                          msgpack::Type::String))
     return false;
   if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_size", true))
     return false;
   if (!verifyIntegerEntry(KernelMap, ".group_segment_fixed_size", true))
     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))
     return false;
   if (!verifyIntegerEntry(KernelMap, ".sgpr_count", true))
     return false;
   if (!verifyIntegerEntry(KernelMap, ".vgpr_count", true))
     return false;

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) {
+    KernelCodeProperties |= amdhsa::KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK;
+  }
+
   return KernelCodeProperties;
 }
 
 amdhsa::kernel_descriptor_t AMDGPUAsmPrinter::getAmdhsaKernelDescriptor(
     const MachineFunction &MF,
     const SIProgramInfo &PI) const {
   const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
   const Function &F = MF.getFunction();

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
 
   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);
+  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/AsmParser/AMDGPUAsmParser.cpp

if (ID == ".amdhsa_group_segment_fixed_size") {
         ImpliedUserSGPRCount += 1;
     } else if (ID == ".amdhsa_wavefront_size32") {
       if (IVersion.Major < 10)
         return Error(IDRange.Start, "directive requires gfx10+", IDRange);
       EnableWavefrontSize32 = Val;
       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,
                      "directive is not supported with architected flat scratch",
                      IDRange);
       PARSE_BITS_ENTRY(KD.compute_pgm_rsrc2,
                        COMPUTE_PGM_RSRC2_ENABLE_PRIVATE_SEGMENT, Val, ValRange);
     } else if (ID == ".amdhsa_enable_private_segment") {

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);
     }
 
+    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.
     ReservedBytes = DE.getBytes(Cursor, 6);

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);
+  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

 ; GCN: .amdhsa_kernel minimal_kernel_inputs
 ; GCN: .amdhsa_user_sgpr_count 15
 ; 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 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
 ; 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
 ; GCN: ; COMPUTE_PGM_RSRC2:USER_SGPR: 15
 define amdgpu_kernel void @minimal_kernel_inputs() {

 ; GCN: .amdhsa_kernel minimal_kernel_inputs
 ; GCN: .amdhsa_user_sgpr_count 15
 ; 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 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
 ; 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
 ; GCN: ; COMPUTE_PGM_RSRC2:USER_SGPR: 15
 define amdgpu_kernel void @minimal_kernel_inputs_with_stack() {

 ; GCN: .amdhsa_kernel queue_ptr
 ; GCN: .amdhsa_user_sgpr_count 15
 ; 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 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
 ; 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
 ; GCN: ; COMPUTE_PGM_RSRC2:USER_SGPR: 15
 define amdgpu_kernel void @queue_ptr() {

 ; GCN: .amdhsa_kernel all_inputs
 ; GCN: .amdhsa_user_sgpr_count 13
 ; 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 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
 ; 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
 ; GCN: ; COMPUTE_PGM_RSRC2:USER_SGPR: 13
 define amdgpu_kernel void @all_inputs() {

llvm/test/CodeGen/AMDGPU/indirect-call-known-callees.ll

 ; CHECK-NEXT: .amdhsa_user_sgpr_count 6
 ; CHECK-NEXT: .amdhsa_user_sgpr_private_segment_buffer 1
 ; CHECK-NEXT: .amdhsa_user_sgpr_dispatch_ptr 0
 ; CHECK-NEXT: .amdhsa_user_sgpr_queue_ptr 0
 ; CHECK-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 0
 ; 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
 ; CHECK-NEXT: .amdhsa_system_sgpr_workgroup_id_z 0
 ; CHECK-NEXT: .amdhsa_system_sgpr_workgroup_info 0
 ; CHECK-NEXT: .amdhsa_system_vgpr_workitem_id 0
 bb:
   %cond = load i1, i1 addrspace(4)* null

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 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
 // 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 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
 // 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 7f000000 00000000
+// OBJDUMP-NEXT: 0070 c1500104 1f0f007f 7f080000 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 5e000000 00000000
+// OBJDUMP-NEXT: 0070 01510104 130f007f 5e080000 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 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
 // 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 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
 // 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/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: 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:
 #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               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:
 # 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

llvm/test/tools/llvm-readobj/ELF/note-amdgpu.test

 # RUN: yaml2obj %s -o %t.o
 # RUN: llvm-readobj --notes %t.o | FileCheck %s --check-prefix=LLVM
 # RUN: llvm-readelf --notes %t.o | FileCheck %s --check-prefix=GNU
 
 # 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:
 # GNU-NEXT:   - .group_segment_fixed_size: 2
 # GNU-NEXT:     .kernarg_segment_align: 4
 # GNU-NEXT:     .kernarg_segment_size: 1
 # GNU-NEXT:     .max_flat_workgroup_size: 8
 # GNU-NEXT:     .name:           foo
 # 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:
 # GNU-NEXT:   - 1
 # GNU-NEXT:   - 0
 # GNU-NEXT: ...
 # GNU-EMPTY:
 # GNU-EMPTY:
 # GNU-NEXT: Displaying notes found in: .note.unknown
 # GNU-NEXT:   Owner                Data size  Description
 # GNU-NEXT:   AMDGPU               0x00000002 Unknown note type: (0x00000101)
 # GNU-NEXT:    description data: ab cd
 # GNU-EMPTY:
 
 # LLVM:      Notes [
 # LLVM-NEXT:   NoteSection {
 # LLVM-NEXT:     Name: .note.foo
 # LLVM-NEXT:     Offset:
 # 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:
 # LLVM-NEXT:   - .group_segment_fixed_size: 2
 # LLVM-NEXT:     .kernarg_segment_align: 4
 # LLVM-NEXT:     .kernarg_segment_size: 1
 # LLVM-NEXT:     .max_flat_workgroup_size: 8
 # LLVM-NEXT:     .name:           foo
 # 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:
 # LLVM-NEXT:   - 1
 # LLVM-NEXT:   - 0
 # LLVM-NEXT: ...
 # LLVM-EMPTY:
 # LLVM-NEXT:     }
 # 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
 # LLVM-NEXT:       Data size: 0x2
 # LLVM-NEXT:       Type: Unknown (0x00000101)
 # LLVM-NEXT:       Description data (
 # LLVM-NEXT:         0000: ABCD                                 |..|
 # LLVM-NEXT:       )

 #     - 1
 #     - 0
 #   amdhsa.kernels:
 #     - .name:   foo
 #       .symbol: foo
 #       .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:
   Class: ELFCLASS64
   Data:  ELFDATA2LSB
   Type:  ET_REL
 Sections:
   - Name:        .note.foo
     Type:        SHT_NOTE
-    Content:     07000000E600000020000000414D44475055000082AE616D646873612E6B65726E656C73918AB92E67726F75705F7365676D656E745F66697865645F73697A6502B62E6B65726E6172675F7365676D656E745F616C69676E04B52E6B65726E6172675F7365676D656E745F73697A6501B82E6D61785F666C61745F776F726B67726F75705F73697A6508A52E6E616D65A3666F6FBB2E707269766174655F7365676D656E745F66697865645F73697A6503AB2E736770725F636F756E7406A72E73796D626F6CA3666F6FAB2E766770725F636F756E7407AF2E7761766566726F6E745F73697A6505AE616D646873612E76657273696F6E9201000000
+    Content:     07000000FB00000020000000414D44475055000082AE616D646873612E6B65726E656C73918BB92E67726F75705F7365676D656E745F66697865645F73697A6502B62E6B65726E6172675F7365676D656E745F616C69676E04B52E6B65726E6172675F7365676D656E745F73697A6501B82E6D61785F666C61745F776F726B67726F75705F73697A6508A52E6E616D65A3666F6FBB2E707269766174655F7365676D656E745F66697865645F73697A6503AB2E736770725F636F756E7406A72E73796D626F6CA3666F6FB32E757365735F64796E616D69635F737461636BC3AB2E766770725F636F756E7407AF2E7761766566726F6E745F73697A6505AE616D646873612E76657273696F6E92010000
   - Name:        .note.unknown
     Type:        SHT_NOTE
     Notes:
       - Name: AMDGPU
         Type: NT_GNU_BUILD_ATTRIBUTE_FUNC
         Desc: 'abcd'