[llvm-branch-commits] [llvm] [SPIRV][SPIRVPrepareGlobals] Map AMD's dynamic LDS 0-element globals to arrays with UINT32_MAX elements (PR #166952)

Juan Manuel Martinez CaamaƱo via llvm-branch-commits llvm-branch-commits at lists.llvm.org
Mon Nov 10 08:08:46 PST 2025


================
@@ -52,6 +76,9 @@ bool SPIRVPrepareGlobals::runOnModule(Module &M) {
   if (GlobalVariable *Bitcode = M.getNamedGlobal("llvm.embedded.module"))
     Changed |= tryExtendLLVMBitcodeMarker(*Bitcode);
 
+  for (GlobalVariable &GV : make_early_inc_range(M.globals()))
+    Changed |= tryExtendDynamicLDSGlobal(GV);
----------------
jmmartinez wrote:

I've added a comment on top of the function:
```cpp
// In HIP, dynamic LDS variables are represented using 0-element global arrays
// in the __shared__ language address-space.
//
//  extern __shared__ int LDS[];
//
// These are not representable in SPIRV directly.
// To represent them, for AMD, we use an array with UINT32_MAX-elements.
// These are reverse translated to 0-element arrays.
bool tryExtendDynamicLDSGlobal(GlobalVariable &GV) {
```

https://github.com/llvm/llvm-project/pull/166952


More information about the llvm-branch-commits mailing list