[amdgpu] Add codegen support for HIP dynamic shared memory.

[hliao]

Summary

• HIP uses an unsized extern array extern __shared__ T s[] to declare the dynamic shared memory, which size is not known at the compile time.

[arsenm]

Needs to handle globalisel too.

I also thought we were trying to get rid of group static size. It's broken with LDS relocations which we need to move towards. Can we just switch directly to using a relocation here?

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

5584

Why does it specifically need to be a 0 sized array? I think this would depend purely on the linkage, or treat any 0 sized type the same way

[scchan]

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

5588

don't you need to make sure whether the static size would give you an offset with the correct alignment?

[hliao]

I just found that change for non-HSA/-PAL environment. I need to check how it works and fit into other tests. So far, that's a critical change to ensure we won't change the original source code too much. Is it possible to address that relocation in a long run (says 1~3 weeks) to avoid the tight schedule.

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

5584

That's the only syntax accepted under HIP clang. At least, zero-sized types should be checked to ensure developers understand the usage of the dynamic shared memory.

[hliao]

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

5588

I remember that size should be always DWORD aligned. Let me check the code calculated that.

[t-tye]

My understanding is this feature is equivalent to the OpenCL dynamic group segment allocation. The runtime would presumably implement it in a similar way.

So the HIP runtime must take the static LDS size, round up to the alignment requirement of the dynamic allocation (OpenCL just uses the maximally aligned OpenCL data type), then add the size of the dynamic LDS. The AQL packet group segment field is set to the total LDS size.

In OpenCL there can be multiple kernel arguments, and the LDS address is passed to each. But for HIP there is only one dynamic area denoted by this weird extern. How is the dynamic LDS storage accessed? Is the address passed as an implicit kernel argument, or does the compiler implicitly use the aligned static LDS size?

[arsenm]

I don't think this actually works since you could have multiple 0 sized objects, and they would both get the same address. I think this has to be an external relocation

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

5584

That's a bit too much HIP specific logic. Also what does this do if there are more than one? How can these return different addresses?

5588

The global has an explicit alignment that needs to be respected

[hliao]

In D82496#2112763, @t-tye wrote:

My understanding is this feature is equivalent to the OpenCL dynamic group segment allocation. The runtime would presumably implement it in a similar way.

So the HIP runtime must take the static LDS size, round up to the alignment requirement of the dynamic allocation (OpenCL just uses the maximally aligned OpenCL data type), then add the size of the dynamic LDS. The AQL packet group segment field is set to the total LDS size.

In OpenCL there can be multiple kernel arguments, and the LDS address is passed to each. But for HIP there is only one dynamic area denoted by this weird extern. How is the dynamic LDS storage accessed? Is the address passed as an implicit kernel argument, or does the compiler implicitly use the aligned static LDS size?

This's the point. To keep compatible with CUDA, multiple dynamically sized arrays in a single kernel must declare a single extern unsized array and uses the address to divide it into multiple arrays by developers themselves. It's in fact quite similar to the local memory parameter in OpenCL. Thus, all extern unsized __shared__ arrays are mapped onto the same address. The developer has the responsibility to divide it into multiple ones. That's in fact is same across HCC and HIP-Clang except that we want to maximize the compatibility and avoid changing source code too much.

[hliao]

Address the alignment issue.

[arsenm]

Also missing the globalisel handling

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.h

31

I don't think this should be mutable

64

Rename to getAllocatedLDSSize()? I think there should be a separate method to get the size plus the roundup to the dynamic alignment

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.h

34

This doesn't need to be MaybeAlign, just Align. Also expand to Dynamic?

Also should probably elaborate that this is used for the case where a dynamically sized global is used

90

Set is the wrong word here; ensureDynamicLDSAlign()?

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

5585

This should check if the allocated size of GV->getValueType() is 0, not special case 0 sized arrays

5589

Should be the alignment of the aggregate itself, not the element type

[arsenm]

For the globalisel part, you'll need D84638 and the global lowering should introduce the intrinsic, not the machine pseudo

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.h

34

This also needs to be added to the MachineFunctionInfo serialization

65

I think having this here actually breaks the calculation of the total size for all of the statically known globals

[hliao]

Add GlobalISel and MIR support.

[hliao]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.h

64

To report the LDS usage accurately, the final LDSSize need to count the padding due to dynamic shared memory alignment. Now, the re-alignment on LDS is explicitly done before the final instruction selection.

65

The re-alignment is done explicitly just before final instruction selection and just once. BTW, getLDSSize should be only valid after instruction selection.

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

2808–2811 ↗

(On Diff #284255)

I think these should remain distinct queries/fields, not fixed up at an arbitrary point. GlobalISel will miss this for example. The asm printer would query the kind that accounts for the dynamic padding

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

2255

This should not special case 0 sized arrays and should check the allocated size

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.h

38

This can still just be Align

92

This isn't a set since it does more than set the value. ensureDynamicLDSAlign?

98–99

This shouldn't be a mutation, but return the aligned up size.

totalLDSAllocSize()?

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

5585

This should not special case 0 sized arrays. This is 0 allocation size

llvm/test/CodeGen/MIR/AMDGPU/machine-function-info.ll

50

The default should be 1

[hliao]

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

2808–2811 ↗

(On Diff #284255)

GlobalISel calls adjustLDSSizeForDynLDSAlign similarly before the finalization of ISel.

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

2255

I tend to be restrictive here to follow how that is used in HIP, zero-sized array always has zero allocated size. If other clients need similar usage but general zero-allocated type, we may enhance accordingly.

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.h

98–99

don't we need to report the accurate usage of LDS? Does that alignment padding need counting as well for the final LDSSize?

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

2808–2811 ↗

(On Diff #284255)

Having extra state that needs to be made consistent is bad. It's better to just track the two independent fields and do the roundup when needed at the end

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

2255

Types don't mean anything. Any 0 sized globals are getting allocated to the same address. We're just going to miscompile other 0 sized types. We have no reason to treat other 0 sized types differently

2257–2258

Should use the aggregate alignment, not the element

[hliao]

Add preFinalizeLowering so that both DAGISel and GISel shares the same path
to adjust LDS size.

[hliao]

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

2257–2258

Propose to add preFinalizeLowering before pseudo instruction expansion so that both GISel and DAGISel have the chance to adjust LDS size.

[arsenm]

llvm/include/llvm/CodeGen/TargetLowering.h

2794

The last thing we need is more callbacks called at random points

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

2257–2258

It's really cleaner to just have this compute what you care about at the point you care about it. Having a point where this needs to be made consistent is both worse from a serialization perspective, and from an optimization point since theoretically we could spill into the padding later

[hliao]

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

2257–2258

But, as we discuss this is a short-term solution before the linker could perform the per-kernel LDS resolution. As the spilling to LDS is not implemented yet, this short-term solution should be kept as simple as possible and finally reverted. To pad LDS for the shared memory array, we have to wait until all static LDS ones are allocated. That's should be the point where amdgcn_groupstatcisize is about to be expanded, i.e. before finalizing lowering.

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

2257–2258

Short term workaround or not, the less mutable state the better. I see no advantage to accumulating the allocated + padding size in a single variable vs. keeping the two separate. You have to track both in MFI anyway, and accumulating like this loses information.

[hliao]

Remove adjustLDSSizeForDynLDSAlign.

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.h

41

Leftover comment: Just before the final selection, LDSSize is adjusted accordingly.

42

This doesn't need to be MaybeAlign, it can just be Align and default to 1

98–99

ensureDynamicLDSAlign(), and don't need to conditionally set it. This also does not need to mutate LDSSize

llvm/test/CodeGen/MIR/AMDGPU/machine-function-info-no-ir.mir

126

Should be 1

[hliao]

Revise following comments.

[hliao]

Rebase

[hliao]

Minor coding style fix.

[arsenm]

llvm/include/llvm/CodeGen/MIRYamlMapping.h

168–171 ↗

(On Diff #286555)

Should add parser tests for these cases

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

59 ↗

(On Diff #286555)

This is an independent field and should not be changed here

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.h

95

This should still not be modifying LDSSize. This is still missing an independent query to give the static + rounded size

[hliao]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

59 ↗

(On Diff #286555)

As the sequence of static LDS allocation and dynamic LDS alignment updates are processed in the program order or reverse of that order, we need to collect all static LDS usage and dynamic LDS alignment. As we remove the previous the one single point adjustment, we need to update LDSSizze if there's any static LDS allocation or dynamic LDS alignment updates.

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.h

95

Valid LDS queries should be done after instruction selection. LDSSize is ONLY modified within instruction selection through static LDS allocation and dynamic LDS alignment update.

[hliao]

llvm/include/llvm/CodeGen/MIRYamlMapping.h

168–171 ↗

(On Diff #286555)

llvm/test/CodeGen/MIR/AMDGPU/machine-function-info-no-ir.mir covers that.

[arsenm]

llvm/include/llvm/CodeGen/MIRYamlMapping.h

168–171 ↗

(On Diff #286555)

It doesn't cover the error cases

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

59 ↗

(On Diff #286555)

OK yes, this is in the right place now. However,r it should be where the alignment is updated

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

5588

This logic should be moved into allocateLDSGlobal

[hliao]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

59 ↗

(On Diff #286555)

As we don't know which one will be processed last, we need to update LDSSize in both cases to ensure the correct one is calculated.

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

5588

the allocation of dynamic LDS is not handled by the compiler. We only collect the alignment.

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

59 ↗

(On Diff #286555)

This is actually terrible and we're going to burn extra padding, but I guess it's conservatively correct

[hliao]

Add the invalid case for alignment parsing.

[hliao]

llvm/include/llvm/CodeGen/MIRYamlMapping.h

168–171 ↗

(On Diff #286555)

Add an invalid case.

[hliao]

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.cpp

59 ↗

(On Diff #286555)

By maintaining StaticLDSSize, the padding for dynamic LDS is done only once. However, we need to continuously update it if there's any static LDS allocation or dynamic LDS alignment updates.

[hliao]

Rebase.

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

2255

Should get global's alignment, not just the type

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

5588

This should take the alignment from the global, not just the type alignment

llvm/test/CodeGen/AMDGPU/GlobalISel/hip.extern.shared.array.ll

11

Needs some tests with larger explicit alignments

llvm/test/CodeGen/MIR/AMDGPU/machine-function-info-dynlds-align-invalid-case.mir

1–2 ↗

(On Diff #286819)

This isn't XFAIL, it's run with not and check the error message output

[hliao]

Check the explicit alignment if any and fix the negative test case.

[hliao]

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

2255

Good catch, fixed in the latest revision.

[arsenm]

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

5585

Should probably add a comment explaining this is dynamically allocated or something

llvm/test/CodeGen/AMDGPU/hip.extern.shared.array.ll

4–11

None of these use an unnaturally high alignment

[hliao]

Add comment and revist the test case.

llvm/include/llvm/CodeGen/TargetLowering.h

void setLibcallCallingConv(RTLIB::Libcall Call, CallingConv::ID CC) {
     LibcallCallingConvs[Call] = CC;
   }
 
   /// Get the CallingConv that should be used for the specified libcall.
   CallingConv::ID getLibcallCallingConv(RTLIB::Libcall Call) const {
     return LibcallCallingConvs[Call];
   }
 
+  /// Execute target specific actions before finalizing target lowering. This
+  /// is used to collect/adjust per-function resource usage before the
+  /// expansion of pseudo instructions.
+  virtual void preFinalizeLowering(MachineFunction &MF) const;
+
   /// Execute target specific actions to finalize target lowering.
   /// This is used to set extra flags in MachineFrameInformation and freezing
   /// the set of reserved registers.
   /// The default implementation just freezes the set of reserved registers.
   virtual void finalizeLowering(MachineFunction &MF) const;
 
   //===----------------------------------------------------------------------===//
   //  GlobalISel Hooks

llvm/lib/CodeGen/FinalizeISel.cpp

 char &llvm::FinalizeISelID = FinalizeISel::ID;
 INITIALIZE_PASS(FinalizeISel, DEBUG_TYPE,
                 "Finalize ISel and expand pseudo-instructions", false, false)
 
 bool FinalizeISel::runOnMachineFunction(MachineFunction &MF) {
   bool Changed = false;
   const TargetLowering *TLI = MF.getSubtarget().getTargetLowering();
 
+  TLI->preFinalizeLowering(MF);
+
   // Iterate through each instruction in the function, looking for pseudos.
   for (MachineFunction::iterator I = MF.begin(), E = MF.end(); I != E; ++I) {
     MachineBasicBlock *MBB = &*I;
     for (MachineBasicBlock::iterator MBBI = MBB->begin(), MBBE = MBB->end();
          MBBI != MBBE; ) {
       MachineInstr &MI = *MBBI++;
 
       // If MI is a pseudo, expand it.

llvm/lib/CodeGen/TargetLoweringBase.cpp

int TargetLoweringBase::getSqrtRefinementSteps(EVT VT,
   return getOpRefinementSteps(true, VT, getRecipEstimateForFunc(MF));
 }
 
 int TargetLoweringBase::getDivRefinementSteps(EVT VT,
                                               MachineFunction &MF) const {
   return getOpRefinementSteps(false, VT, getRecipEstimateForFunc(MF));
 }
 
+void TargetLoweringBase::preFinalizeLowering(MachineFunction &MF) const {
+  // DO NOTHING.
+}
+
 void TargetLoweringBase::finalizeLowering(MachineFunction &MF) const {
   MF.getRegInfo().freezeReservedRegs(MF);
 }
 
 MachineMemOperand::Flags
 TargetLoweringBase::getLoadMemOperandFlags(const LoadInst &LI,
                                            const DataLayout &DL) const {
   MachineMemOperand::Flags Flags = MachineMemOperand::MOLoad;

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

if (AS == AMDGPUAS::LOCAL_ADDRESS || AS == AMDGPUAS::REGION_ADDRESS) {
     // TODO: We could emit code to handle the initialization somewhere.
     if (!AMDGPUTargetLowering::hasDefinedInitializer(GV)) {
       const SITargetLowering *TLI = ST.getTargetLowering();
       if (!TLI->shouldUseLDSConstAddress(GV)) {
         MI.getOperand(1).setTargetFlags(SIInstrInfo::MO_ABS32_LO);
         return true; // Leave in place;
       }
 
+      if (AS == AMDGPUAS::LOCAL_ADDRESS && GV->hasExternalLinkage()) {
+        Type *Ty = GV->getValueType();
+        if (B.getDataLayout().getTypeAllocSize(Ty).isZero()) {
+          // Adjust alignment for that dynamic shared memory array.
+          MFI->setDynLDSAlign(B.getDataLayout().getABITypeAlign(Ty));
+          LLT S32 = LLT::scalar(32);
+          auto Sz =
+              B.buildIntrinsic(Intrinsic::amdgcn_groupstaticsize, {S32}, false);
+          B.buildIntToPtr(DstReg, Sz);
+          MI.eraseFromParent();
+          return true;
+        }
+      }
+
       B.buildConstant(
           DstReg,
           MFI->allocateLDSGlobal(B.getDataLayout(), *cast<GlobalVariable>(GV)));
       MI.eraseFromParent();
       return true;
     }
 
     const Function &Fn = MF.getFunction();

llvm/lib/Target/AMDGPU/AMDGPUMachineFunction.h

 //===-- AMDGPUMachineFunctionInfo.h -------------------------------*- C++ -*-=//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPUMACHINEFUNCTION_H
 #define LLVM_LIB_TARGET_AMDGPU_AMDGPUMACHINEFUNCTION_H
 
+#include "Utils/AMDGPUBaseInfo.h"
 #include "llvm/ADT/DenseMap.h"
 #include "llvm/CodeGen/MachineFunction.h"
-#include "Utils/AMDGPUBaseInfo.h"
+#include "llvm/Support/Alignment.h"
 
 namespace llvm {
 
 class GCNSubtarget;
 
 class AMDGPUMachineFunction : public MachineFunctionInfo {
   /// A map to keep track of local memory objects and their offsets within the
   /// local memory space.
   SmallDenseMap<const GlobalValue *, unsigned, 4> LocalMemoryObjects;
 
 protected:
   uint64_t ExplicitKernArgSize = 0; // Cache for this.
   Align MaxKernArgAlign;        // Cache for this.
 
   /// Number of bytes in the LDS that are being used.
   unsigned LDSSize = 0;
 
+  /// Align for dynamic shared memory if any. Dynamic shared memory is
+  /// allocated directly after the static one, i.e., LDSSize. Need to pad
+  /// LDSSize to ensure that dynamic one is aligned accordingly.
+  /// The maximal alignment is updated during IR translation or lowering
+  /// stages. Just before the final selection, LDSSize is adjusted accordingly.
+  MaybeAlign DynLDSAlign;
+
   // State of MODE register, assumed FP mode.
   AMDGPU::SIModeRegisterDefaults Mode;
 
   // Kernels + shaders. i.e. functions called by the driver and not called
   // by other functions.
   bool IsEntryFunction = false;
 
   bool NoSignedZerosFPMath = false;
 
   // Function may be memory bound.
   bool MemoryBound = false;
 
   // Kernel may need limited waves per EU for better performance.
   bool WaveLimiter = false;
 
 public:
   AMDGPUMachineFunction(const MachineFunction &MF);
 
   uint64_t getExplicitKernArgSize() const {
     return ExplicitKernArgSize;
   }
 
   unsigned getMaxKernArgAlign() const { return MaxKernArgAlign.value(); }
 
   unsigned getLDSSize() const {
     return LDSSize;
   }
 
   AMDGPU::SIModeRegisterDefaults getMode() const {
     return Mode;
   }
 
   bool isEntryFunction() const {
     return IsEntryFunction;
   }
 
   bool hasNoSignedZerosFPMath() const {
     return NoSignedZerosFPMath;
   }
 
   bool isMemoryBound() const {
     return MemoryBound;
   }
 
   bool needsWaveLimiter() const {
     return WaveLimiter;
   }
 
   unsigned allocateLDSGlobal(const DataLayout &DL, const GlobalVariable &GV);
+
+  MaybeAlign getDynLDSAlign() const { return DynLDSAlign; }
+
+  void setDynLDSAlign(Align A) {
+    if (!DynLDSAlign || A > *DynLDSAlign)
+      DynLDSAlign = A;
+  }
+
+  void adjustLDSSizeForDynLDSAlign() {
+    if (!DynLDSAlign)
+      return;
+    LDSSize = alignTo(LDSSize, *DynLDSAlign);
+  }
 };
 
 }
 #endif

llvm/lib/Target/AMDGPU/SIISelLowering.h

bool checkAsmConstraintVal(SDValue Op,
                              const std::string &Constraint,
                              uint64_t Val) const;
   bool checkAsmConstraintValA(SDValue Op,
                               uint64_t Val,
                               unsigned MaxSize = 64) const;
   SDValue copyToM0(SelectionDAG &DAG, SDValue Chain, const SDLoc &DL,
                    SDValue V) const;
 
+  void preFinalizeLowering(MachineFunction &MF) const override;
+
   void finalizeLowering(MachineFunction &MF) const override;
 
   void computeKnownBitsForFrameIndex(int FrameIdx,
                                      KnownBits &Known,
                                      const MachineFunction &MF) const override;
 
   Align computeKnownAlignForTargetInstr(GISelKnownBits &Analysis, Register R,
                                         const MachineRegisterInfo &MRI,

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

buildPCRelGlobalAddress(SelectionDAG &DAG, const GlobalValue *GV,
   }
   return DAG.getNode(AMDGPUISD::PC_ADD_REL_OFFSET, DL, PtrVT, PtrLo, PtrHi);
 }
 
 SDValue SITargetLowering::LowerGlobalAddress(AMDGPUMachineFunction *MFI,
                                              SDValue Op,
                                              SelectionDAG &DAG) const {
   GlobalAddressSDNode *GSD = cast<GlobalAddressSDNode>(Op);
+  SDLoc DL(GSD);
+  EVT PtrVT = Op.getValueType();
+
   const GlobalValue *GV = GSD->getGlobal();
   if ((GSD->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS &&
        shouldUseLDSConstAddress(GV)) ||
       GSD->getAddressSpace() == AMDGPUAS::REGION_ADDRESS ||
-      GSD->getAddressSpace() == AMDGPUAS::PRIVATE_ADDRESS)
+      GSD->getAddressSpace() == AMDGPUAS::PRIVATE_ADDRESS) {
+    if (GSD->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS &&
+        GV->hasExternalLinkage()) {
+      Type *Ty = GV->getValueType();
+      if (DAG.getDataLayout().getTypeAllocSize(Ty).isZero()) {
+        assert(PtrVT == MVT::i32 && "32-bit pointer is expected.");
+        // Adjust alignment for that dynamic shared memory array.
+        MFI->setDynLDSAlign(DAG.getDataLayout().getABITypeAlign(Ty));
+        return SDValue(
+            DAG.getMachineNode(AMDGPU::GET_GROUPSTATICSIZE, DL, PtrVT), 0);
+      }
+    }
     return AMDGPUTargetLowering::LowerGlobalAddress(MFI, Op, DAG);
-
-  SDLoc DL(GSD);
-  EVT PtrVT = Op.getValueType();
+  }
 
   if (GSD->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
     SDValue GA = DAG.getTargetGlobalAddress(GV, DL, MVT::i32, GSD->getOffset(),
                                             SIInstrInfo::MO_ABS32_LO);
     return DAG.getNode(AMDGPUISD::LDS, DL, MVT::i32, GA);
   }
 
   if (shouldEmitFixup(GV))
bool SITargetLowering::checkAsmConstraintValA(SDValue Op,
   if ((Size == 16 && AMDGPU::isInlinableLiteral16(Val, HasInv2Pi)) ||
       (Size == 32 && AMDGPU::isInlinableLiteral32(Val, HasInv2Pi)) ||
       (Size == 64 && AMDGPU::isInlinableLiteral64(Val, HasInv2Pi))) {
     return true;
   }
   return false;
 }
 
+void SITargetLowering::preFinalizeLowering(MachineFunction &MF) const {
+  SIMachineFunctionInfo *Info = MF.getInfo<SIMachineFunctionInfo>();
+  // Adjust LDS to align the dynamic shared array.
+  Info->adjustLDSSizeForDynLDSAlign();
+}
+
 // Figure out which registers should be reserved for stack access. Only after
 // the function is legalized do we know all of the non-spill stack objects or if
 // calls are present.
 void SITargetLowering::finalizeLowering(MachineFunction &MF) const {
   MachineRegisterInfo &MRI = MF.getRegInfo();
   SIMachineFunctionInfo *Info = MF.getInfo<SIMachineFunctionInfo>();
   const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
   const SIRegisterInfo *TRI = Subtarget->getRegisterInfo();

llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h

static void mapping(IO &YamlIO, SIMode &Mode) {
     YamlIO.mapOptional("fp64-fp16-output-denormals", Mode.FP64FP16OutputDenormals, true);
   }
 };
 
 struct SIMachineFunctionInfo final : public yaml::MachineFunctionInfo {
   uint64_t ExplicitKernArgSize = 0;
   unsigned MaxKernArgAlign = 0;
   unsigned LDSSize = 0;
+  MaybeAlign DynLDSAlign;
   bool IsEntryFunction = false;
   bool NoSignedZerosFPMath = false;
   bool MemoryBound = false;
   bool WaveLimiter = false;
   bool HasSpilledSGPRs = false;
   bool HasSpilledVGPRs = false;
   uint32_t HighBitsOf32BitAddress = 0;
 

 };
 
 template <> struct MappingTraits<SIMachineFunctionInfo> {
   static void mapping(IO &YamlIO, SIMachineFunctionInfo &MFI) {
     YamlIO.mapOptional("explicitKernArgSize", MFI.ExplicitKernArgSize,
                        UINT64_C(0));
     YamlIO.mapOptional("maxKernArgAlign", MFI.MaxKernArgAlign, 0u);
     YamlIO.mapOptional("ldsSize", MFI.LDSSize, 0u);
+    YamlIO.mapOptional("dynLDSAlign", MFI.DynLDSAlign, MaybeAlign());
     YamlIO.mapOptional("isEntryFunction", MFI.IsEntryFunction, false);
     YamlIO.mapOptional("noSignedZerosFPMath", MFI.NoSignedZerosFPMath, false);
     YamlIO.mapOptional("memoryBound", MFI.MemoryBound, false);
     YamlIO.mapOptional("waveLimiter", MFI.WaveLimiter, false);
     YamlIO.mapOptional("hasSpilledSGPRs", MFI.HasSpilledSGPRs, false);
     YamlIO.mapOptional("hasSpilledVGPRs", MFI.HasSpilledVGPRs, false);
     YamlIO.mapOptional("scratchRSrcReg", MFI.ScratchRSrcReg,
                        StringValue("$private_rsrc_reg"));

llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp

convertArgumentInfo(const AMDGPUFunctionArgInfo &ArgInfo,
 
   if (Any)
     return AI;
 
   return None;
 }
 
 yaml::SIMachineFunctionInfo::SIMachineFunctionInfo(
-  const llvm::SIMachineFunctionInfo& MFI,
-  const TargetRegisterInfo &TRI)
+    const llvm::SIMachineFunctionInfo &MFI, const TargetRegisterInfo &TRI)
     : ExplicitKernArgSize(MFI.getExplicitKernArgSize()),
-    MaxKernArgAlign(MFI.getMaxKernArgAlign()),
-    LDSSize(MFI.getLDSSize()),
-    IsEntryFunction(MFI.isEntryFunction()),
+      MaxKernArgAlign(MFI.getMaxKernArgAlign()), LDSSize(MFI.getLDSSize()),
+      DynLDSAlign(MFI.getDynLDSAlign()), IsEntryFunction(MFI.isEntryFunction()),
       NoSignedZerosFPMath(MFI.hasNoSignedZerosFPMath()),
-    MemoryBound(MFI.isMemoryBound()),
-    WaveLimiter(MFI.needsWaveLimiter()),
+      MemoryBound(MFI.isMemoryBound()), WaveLimiter(MFI.needsWaveLimiter()),
       HasSpilledSGPRs(MFI.hasSpilledSGPRs()),
       HasSpilledVGPRs(MFI.hasSpilledVGPRs()),
       HighBitsOf32BitAddress(MFI.get32BitAddressHighBits()),
       ScratchRSrcReg(regToString(MFI.getScratchRSrcReg(), TRI)),
       FrameOffsetReg(regToString(MFI.getFrameOffsetReg(), TRI)),
       StackPtrOffsetReg(regToString(MFI.getStackPtrOffsetReg(), TRI)),
-    ArgInfo(convertArgumentInfo(MFI.getArgInfo(), TRI)),
-    Mode(MFI.getMode()) {}
+      ArgInfo(convertArgumentInfo(MFI.getArgInfo(), TRI)), Mode(MFI.getMode()) {
+}
 
 void yaml::SIMachineFunctionInfo::mappingImpl(yaml::IO &YamlIO) {
   MappingTraits<SIMachineFunctionInfo>::mapping(YamlIO, *this);
 }
 
 bool SIMachineFunctionInfo::initializeBaseYamlFields(
   const yaml::SIMachineFunctionInfo &YamlMFI) {
   ExplicitKernArgSize = YamlMFI.ExplicitKernArgSize;
   MaxKernArgAlign = assumeAligned(YamlMFI.MaxKernArgAlign);
   LDSSize = YamlMFI.LDSSize;
+  DynLDSAlign = YamlMFI.DynLDSAlign;
   HighBitsOf32BitAddress = YamlMFI.HighBitsOf32BitAddress;
   IsEntryFunction = YamlMFI.IsEntryFunction;
   NoSignedZerosFPMath = YamlMFI.NoSignedZerosFPMath;
   MemoryBound = YamlMFI.MemoryBound;
   WaveLimiter = YamlMFI.WaveLimiter;
   HasSpilledSGPRs = YamlMFI.HasSpilledSGPRs;
   HasSpilledVGPRs = YamlMFI.HasSpilledVGPRs;
   return false;

llvm/test/CodeGen/AMDGPU/GlobalISel/hip.extern.shared.array.ll

+; RUN: llc -global-isel -mtriple=amdgcn--amdhsa -mcpu=gfx900 -verify-machineinstrs -o - %s | FileCheck %s
+
+@lds0 = addrspace(3) global [512 x float] undef, align 4
+@lds1 = addrspace(3) global [256 x float] undef, align 4
+@lds2 = addrspace(3) global [4096 x float] undef, align 4
+@lds3 = addrspace(3) global [67 x i8] undef, align 4
+
+@dynamic_shared0 = external addrspace(3) global [0 x float], align 4
+@dynamic_shared1 = external addrspace(3) global [0 x double], align 4
+
+; CHECK-LABEL: {{^}}dynamic_shared_array_0:
+; CHECK: v_add_u32_e32 v{{[0-9]+}}, 0x800, v{{[0-9]+}}
+define amdgpu_kernel void @dynamic_shared_array_0(float addrspace(1)* %out) {
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %arrayidx0 = getelementptr inbounds [512 x float], [512 x float] addrspace(3)* @lds0, i32 0, i32 %tid.x
+  %val0 = load float, float addrspace(3)* %arrayidx0, align 4
+  %arrayidx1 = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared0, i32 0, i32 %tid.x
+  store float %val0, float addrspace(3)* %arrayidx1, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}dynamic_shared_array_1:
+; CHECK: v_lshlrev_b32_e32 {{v[0-9]+}}, 2, {{v[0-9]+}}
+; CHECK: v_lshlrev_b32_e32 {{v[0-9]+}}, 2, {{v[0-9]+}}
+; CHECK: v_lshlrev_b32_e32 [[IDX:v[0-9]+]], 2, {{v[0-9]+}}
+; CHECK: v_add_u32_e32 {{v[0-9]+}}, 0xc00, [[IDX]]
+define amdgpu_kernel void @dynamic_shared_array_1(float addrspace(1)* %out, i32 %cond) {
+entry:
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %idx.0 = add nsw i32 %tid.x, 64
+  %tmp = icmp eq i32 %cond, 0
+  br i1 %tmp, label %if, label %else
+
+if:                                               ; preds = %entry
+  %arrayidx0 = getelementptr inbounds [512 x float], [512 x float] addrspace(3)* @lds0, i32 0, i32 %idx.0
+  %val0 = load float, float addrspace(3)* %arrayidx0, align 4
+  br label %endif
+
+else:                                             ; preds = %entry
+  %arrayidx1 = getelementptr inbounds [256 x float], [256 x float] addrspace(3)* @lds1, i32 0, i32 %idx.0
+  %val1 = load float, float addrspace(3)* %arrayidx1, align 4
+  br label %endif
+
+endif:                                            ; preds = %else, %if
+  %val = phi float [ %val0, %if ], [ %val1, %else ]
+  %arrayidx = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared0, i32 0, i32 %tid.x
+  store float %val, float addrspace(3)* %arrayidx, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}dynamic_shared_array_2:
+; CHECK: v_lshlrev_b32_e32 [[IDX:v[0-9]+]], 2, {{v[0-9]+}}
+; CHECK: v_add_u32_e32 {{v[0-9]+}}, 0x4000, [[IDX]]
+define amdgpu_kernel void @dynamic_shared_array_2(i32 %idx) {
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %vidx = add i32 %tid.x, %idx
+  %arrayidx0 = getelementptr inbounds [4096 x float], [4096 x float] addrspace(3)* @lds2, i32 0, i32 %vidx
+  %val0 = load float, float addrspace(3)* %arrayidx0, align 4
+  %arrayidx1 = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared0, i32 0, i32 %tid.x
+  store float %val0, float addrspace(3)* %arrayidx1, align 4
+  ret void
+}
+
+; The offset to the dynamic shared memory array should be aligned on the type
+; specified.
+; CHECK-LABEL: {{^}}dynamic_shared_array_3:
+; CHECK: v_lshlrev_b32_e32 [[IDX:v[0-9]+]], 2, {{v[0-9]+}}
+; CHECK: v_add_u32_e32 {{v[0-9]+}}, 0x44, [[IDX]]
+define amdgpu_kernel void @dynamic_shared_array_3(i32 %idx) {
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %vidx = add i32 %tid.x, %idx
+  %arrayidx0 = getelementptr inbounds [67 x i8], [67 x i8] addrspace(3)* @lds3, i32 0, i32 %vidx
+  %val0 = load i8, i8 addrspace(3)* %arrayidx0, align 4
+  %val1 = uitofp i8 %val0 to float
+  %arrayidx1 = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared0, i32 0, i32 %tid.x
+  store float %val1, float addrspace(3)* %arrayidx1, align 4
+  ret void
+}
+
+; The offset to the dynamic shared memory array should be aligned on the
+; maximal one.
+; CHECK-LABEL: {{^}}dynamic_shared_array_4:
+; CHECK: v_mov_b32_e32 [[DYNLDS:v[0-9]+]], 0x48
+; CHECK: v_lshlrev_b32_e32 [[IDX:v[0-9]+]], 2, {{v[0-9]+}}
+; CHECK: v_add_u32_e32 {{v[0-9]+}}, [[DYNLDS]], [[IDX]]
+define amdgpu_kernel void @dynamic_shared_array_4(i32 %idx) {
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %vidx = add i32 %tid.x, %idx
+  %arrayidx0 = getelementptr inbounds [67 x i8], [67 x i8] addrspace(3)* @lds3, i32 0, i32 %vidx
+  %val0 = load i8, i8 addrspace(3)* %arrayidx0, align 4
+  %val1 = uitofp i8 %val0 to float
+  %val2 = uitofp i8 %val0 to double
+  %arrayidx1 = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared0, i32 0, i32 %tid.x
+  store float %val1, float addrspace(3)* %arrayidx1, align 4
+  %arrayidx2 = getelementptr inbounds [0 x double], [0 x double] addrspace(3)* @dynamic_shared1, i32 0, i32 %tid.x
+  store double %val2, double addrspace(3)* %arrayidx2, align 4
+  ret void
+}
+
+declare i32 @llvm.amdgcn.workitem.id.x()

llvm/test/CodeGen/AMDGPU/hip.extern.shared.array.ll

+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx900 -verify-machineinstrs -o - %s | FileCheck %s
+
+@lds0 = addrspace(3) global [512 x float] undef, align 4
+@lds1 = addrspace(3) global [256 x float] undef, align 4
+@lds2 = addrspace(3) global [4096 x float] undef, align 4
+@lds3 = addrspace(3) global [67 x i8] undef, align 4
+
+@dynamic_shared0 = external addrspace(3) global [0 x float], align 4
+@dynamic_shared1 = external addrspace(3) global [0 x double], align 4
+
+; CHECK-LABEL: {{^}}dynamic_shared_array_0:
+; CHECK: v_add_u32_e32 v{{[0-9]+}}, 0x800, v{{[0-9]+}}
+define amdgpu_kernel void @dynamic_shared_array_0(float addrspace(1)* %out) {
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %arrayidx0 = getelementptr inbounds [512 x float], [512 x float] addrspace(3)* @lds0, i32 0, i32 %tid.x
+  %val0 = load float, float addrspace(3)* %arrayidx0, align 4
+  %arrayidx1 = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared0, i32 0, i32 %tid.x
+  store float %val0, float addrspace(3)* %arrayidx1, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}dynamic_shared_array_1:
+; CHECK: v_mov_b32_e32 [[DYNLDS:v[0-9]+]], 0xc00
+; CHECK: v_lshl_add_u32 {{v[0-9]+}}, {{v[0-9]+}}, 2, [[DYNLDS]]
+define amdgpu_kernel void @dynamic_shared_array_1(float addrspace(1)* %out, i32 %cond) {
+entry:
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %idx.0 = add nsw i32 %tid.x, 64
+  %tmp = icmp eq i32 %cond, 0
+  br i1 %tmp, label %if, label %else
+
+if:                                               ; preds = %entry
+  %arrayidx0 = getelementptr inbounds [512 x float], [512 x float] addrspace(3)* @lds0, i32 0, i32 %idx.0
+  %val0 = load float, float addrspace(3)* %arrayidx0, align 4
+  br label %endif
+
+else:                                             ; preds = %entry
+  %arrayidx1 = getelementptr inbounds [256 x float], [256 x float] addrspace(3)* @lds1, i32 0, i32 %idx.0
+  %val1 = load float, float addrspace(3)* %arrayidx1, align 4
+  br label %endif
+
+endif:                                            ; preds = %else, %if
+  %val = phi float [ %val0, %if ], [ %val1, %else ]
+  %arrayidx = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared0, i32 0, i32 %tid.x
+  store float %val, float addrspace(3)* %arrayidx, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}dynamic_shared_array_2:
+; CHECK: v_mov_b32_e32 [[DYNLDS:v[0-9]+]], 0x4000
+; CHECK: v_lshl_add_u32 {{v[0-9]+}}, {{v[0-9]+}}, 2, [[DYNLDS]]
+define amdgpu_kernel void @dynamic_shared_array_2(i32 %idx) {
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %vidx = add i32 %tid.x, %idx
+  %arrayidx0 = getelementptr inbounds [4096 x float], [4096 x float] addrspace(3)* @lds2, i32 0, i32 %vidx
+  %val0 = load float, float addrspace(3)* %arrayidx0, align 4
+  %arrayidx1 = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared0, i32 0, i32 %tid.x
+  store float %val0, float addrspace(3)* %arrayidx1, align 4
+  ret void
+}
+
+; The offset to the dynamic shared memory array should be aligned on the type
+; specified.
+; CHECK-LABEL: {{^}}dynamic_shared_array_3:
+; CHECK: v_mov_b32_e32 [[DYNLDS:v[0-9]+]], 0x44
+; CHECK: v_lshl_add_u32 {{v[0-9]+}}, {{v[0-9]+}}, 2, [[DYNLDS]]
+define amdgpu_kernel void @dynamic_shared_array_3(i32 %idx) {
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %vidx = add i32 %tid.x, %idx
+  %arrayidx0 = getelementptr inbounds [67 x i8], [67 x i8] addrspace(3)* @lds3, i32 0, i32 %vidx
+  %val0 = load i8, i8 addrspace(3)* %arrayidx0, align 4
+  %val1 = uitofp i8 %val0 to float
+  %arrayidx1 = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared0, i32 0, i32 %tid.x
+  store float %val1, float addrspace(3)* %arrayidx1, align 4
+  ret void
+}
+
+; The offset to the dynamic shared memory array should be aligned on the
+; maximal one.
+; CHECK-LABEL: {{^}}dynamic_shared_array_4:
+; CHECK: v_mov_b32_e32 [[DYNLDS:v[0-9]+]], 0x48
+; CHECK-DAG: v_lshl_add_u32 {{v[0-9]+}}, {{v[0-9]+}}, 2, [[DYNLDS]]
+; CHECK-DAG: v_lshl_add_u32 {{v[0-9]+}}, {{v[0-9]+}}, 3, [[DYNLDS]]
+define amdgpu_kernel void @dynamic_shared_array_4(i32 %idx) {
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %vidx = add i32 %tid.x, %idx
+  %arrayidx0 = getelementptr inbounds [67 x i8], [67 x i8] addrspace(3)* @lds3, i32 0, i32 %vidx
+  %val0 = load i8, i8 addrspace(3)* %arrayidx0, align 4
+  %val1 = uitofp i8 %val0 to float
+  %val2 = uitofp i8 %val0 to double
+  %arrayidx1 = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared0, i32 0, i32 %tid.x
+  store float %val1, float addrspace(3)* %arrayidx1, align 4
+  %arrayidx2 = getelementptr inbounds [0 x double], [0 x double] addrspace(3)* @dynamic_shared1, i32 0, i32 %tid.x
+  store double %val2, double addrspace(3)* %arrayidx2, align 4
+  ret void
+}
+
+declare i32 @llvm.amdgcn.workitem.id.x()

llvm/test/CodeGen/MIR/AMDGPU/machine-function-info-no-ir.mir

 # RUN: llc -mtriple=amdgcn-amd-amdhsa -run-pass=none -verify-machineinstrs %s -o - | FileCheck -check-prefixes=FULL,ALL %s
 # RUN: llc -mtriple=amdgcn-amd-amdhsa -run-pass=none -simplify-mir -verify-machineinstrs %s -o - | FileCheck -check-prefixes=SIMPLE,ALL %s
 
 
 ---
 # ALL-LABEL: name: kernel0
 # FULL: machineFunctionInfo:
 # FULL-NEXT: explicitKernArgSize: 128
 # FULL-NEXT: maxKernArgAlign: 64
 # FULL-NEXT: ldsSize: 2048
+# FULL-NEXT: dynLDSAlign: 0
 # FULL-NEXT: isEntryFunction: true
 # FULL-NEXT: noSignedZerosFPMath: false
 # FULL-NEXT: memoryBound:     true
 # FULL-NEXT: waveLimiter:     true
 # FULL-NEXT: hasSpilledSGPRs: false
 # FULL-NEXT: hasSpilledVGPRs: false
 # FULL-NEXT: scratchRSrcReg:  '$sgpr8_sgpr9_sgpr10_sgpr11'
 # FULL-NEXT: frameOffsetReg:  '$sgpr12'

 
 # FIXME: Should be able to not print section for simple
 ---
 # ALL-LABEL: name: no_mfi
 # FULL: machineFunctionInfo:
 # FULL-NEXT: explicitKernArgSize: 0
 # FULL-NEXT: maxKernArgAlign: 1
 # FULL-NEXT: ldsSize: 0
+# FULL-NEXT: dynLDSAlign: 0
 # FULL-NEXT: isEntryFunction: false
 # FULL-NEXT: noSignedZerosFPMath: false
 # FULL-NEXT: memoryBound:     false
 # FULL-NEXT: waveLimiter:     false
 # FULL-NEXT: hasSpilledSGPRs: false
 # FULL-NEXT: hasSpilledVGPRs: false
 # FULL-NEXT: scratchRSrcReg:  '$private_rsrc_reg'
 # FULL-NEXT: frameOffsetReg:  '$fp_reg'

 ...
 
 ---
 # ALL-LABEL: name: empty_mfi
 # FULL: machineFunctionInfo:
 # FULL-NEXT: explicitKernArgSize: 0
 # FULL-NEXT: maxKernArgAlign: 1
 # FULL-NEXT: ldsSize: 0
+# FULL-NEXT: dynLDSAlign: 0
 # FULL-NEXT: isEntryFunction: false
 # FULL-NEXT: noSignedZerosFPMath: false
 # FULL-NEXT: memoryBound:     false
 # FULL-NEXT: waveLimiter:     false
 # FULL-NEXT: hasSpilledSGPRs: false
 # FULL-NEXT: hasSpilledVGPRs: false
 # FULL-NEXT: scratchRSrcReg:  '$private_rsrc_reg'
 # FULL-NEXT: frameOffsetReg:  '$fp_reg'

 ...
 
 ---
 # ALL-LABEL: name: empty_mfi_entry_func
 # FULL: machineFunctionInfo:
 # FULL-NEXT: explicitKernArgSize: 0
 # FULL-NEXT: maxKernArgAlign: 1
 # FULL-NEXT: ldsSize: 0
+# FULL-NEXT: dynLDSAlign: 0
 # FULL-NEXT: isEntryFunction: true
 # FULL-NEXT: noSignedZerosFPMath: false
 # FULL-NEXT: memoryBound:     false
 # FULL-NEXT: waveLimiter:     false
 # FULL-NEXT: hasSpilledSGPRs: false
 # FULL-NEXT: hasSpilledVGPRs: false
 # FULL-NEXT: scratchRSrcReg:  '$private_rsrc_reg'
 # FULL-NEXT: frameOffsetReg:  '$fp_reg'
machineFunctionInfo:
   hasSpilledSGPRs: true
   hasSpilledVGPRs: true
 
 body:             |
   bb.0:
     S_ENDPGM 0
 
 ...
+
+---
+# ALL-LABEL: name: dyn_lds_with_alignment
+
+# FULL: ldsSize: 0
+# FULL-NEXT: dynLDSAlign: 8
+
+# SIMPLE: dynLDSAlign: 8
+name: dyn_lds_with_alignment
+machineFunctionInfo:
+  dynLDSAlign: 8
+
+body:             |
+  bb.0:
+    S_ENDPGM 0
+
+...

llvm/test/CodeGen/MIR/AMDGPU/machine-function-info.ll

 ; RUN: llc -mtriple=amdgcn-mesa-mesa3d -mcpu=tahiti -stop-after finalize-isel -o %t.mir %s
 ; RUN: llc -run-pass=none -verify-machineinstrs %t.mir -o - | FileCheck %s
 
 ; Test that SIMachineFunctionInfo can be round trip serialized through
 ; MIR.
 
 @lds = addrspace(3) global [512 x float] undef, align 4
 
 ; CHECK-LABEL: {{^}}name: kernel
 ; CHECK: machineFunctionInfo:
 ; CHECK-NEXT: explicitKernArgSize: 128
 ; CHECK-NEXT: maxKernArgAlign: 64
 ; CHECK-NEXT: ldsSize: 0
+; CHECK-NEXT: dynLDSAlign: 0
 ; CHECK-NEXT: isEntryFunction: true
 ; CHECK-NEXT: noSignedZerosFPMath: false
 ; CHECK-NEXT: memoryBound: false
 ; CHECK-NEXT: waveLimiter: false
 ; CHECK-NEXT: hasSpilledSGPRs: false
 ; CHECK-NEXT: hasSpilledVGPRs: false
 ; CHECK-NEXT: scratchRSrcReg:  '$sgpr96_sgpr97_sgpr98_sgpr99'
 ; CHECK-NEXT: frameOffsetReg:  '$fp_reg'
define amdgpu_kernel void @kernel(i32 %arg0, i64 %arg1, <16 x i32> %arg2) {
   ret void
 }
 
 ; CHECK-LABEL: {{^}}name: ps_shader
 ; CHECK: machineFunctionInfo:
 ; CHECK-NEXT: explicitKernArgSize: 0
 ; CHECK-NEXT: maxKernArgAlign: 1
 ; CHECK-NEXT: ldsSize: 0
+; CHECK-NEXT: dynLDSAlign: 0
 ; CHECK-NEXT: isEntryFunction: true
 ; CHECK-NEXT: noSignedZerosFPMath: false
 ; CHECK-NEXT: memoryBound: false
 ; CHECK-NEXT: waveLimiter: false
 ; CHECK-NEXT: hasSpilledSGPRs: false
 ; CHECK-NEXT: hasSpilledVGPRs: false
 ; CHECK-NEXT: scratchRSrcReg:  '$sgpr96_sgpr97_sgpr98_sgpr99'
 ; CHECK-NEXT: frameOffsetReg:  '$fp_reg'
define amdgpu_ps void @ps_shader(i32 %arg0, i32 inreg %arg1) {
   ret void
 }
 
 ; CHECK-LABEL: {{^}}name: function
 ; CHECK: machineFunctionInfo:
 ; CHECK-NEXT: explicitKernArgSize: 0
 ; CHECK-NEXT: maxKernArgAlign: 1
 ; CHECK-NEXT: ldsSize: 0
+; CHECK-NEXT: dynLDSAlign: 0
 ; CHECK-NEXT: isEntryFunction: false
 ; CHECK-NEXT: noSignedZerosFPMath: false
 ; CHECK-NEXT: memoryBound: false
 ; CHECK-NEXT: waveLimiter: false
 ; CHECK-NEXT: hasSpilledSGPRs: false
 ; CHECK-NEXT: hasSpilledVGPRs: false
 ; CHECK-NEXT: scratchRSrcReg: '$sgpr0_sgpr1_sgpr2_sgpr3'
 ; CHECK-NEXT: frameOffsetReg: '$sgpr33'
define void @function() {
   ret void
 }
 
 ; CHECK-LABEL: {{^}}name: function_nsz
 ; CHECK: machineFunctionInfo:
 ; CHECK-NEXT: explicitKernArgSize: 0
 ; CHECK-NEXT: maxKernArgAlign: 1
 ; CHECK-NEXT: ldsSize: 0
+; CHECK-NEXT: dynLDSAlign: 0
 ; CHECK-NEXT: isEntryFunction: false
 ; CHECK-NEXT: noSignedZerosFPMath: true
 ; CHECK-NEXT: memoryBound: false
 ; CHECK-NEXT: waveLimiter: false
 ; CHECK-NEXT: hasSpilledSGPRs: false
 ; CHECK-NEXT: hasSpilledVGPRs: false
 ; CHECK-NEXT: scratchRSrcReg: '$sgpr0_sgpr1_sgpr2_sgpr3'
 ; CHECK-NEXT: frameOffsetReg: '$sgpr33'