[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]

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

In HIP, dynamic LDS variables are represented using `0-element` global arrays in the `__shared__` language address-space.

```cpp
  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 later in AMD's SPIRV runtime pipeline (in [SPIRVReader.cpp](https://github.com/ROCm/SPIRV-LLVM-Translator/blob/8cb74e264ddcde89f62354544803dc8cdbac148d/lib/SPIRV/SPIRVReader.cpp#L358)).

Stacked over https://github.com/llvm/llvm-project/pull/166950

>From df43b6346491e7b78ff27956da4bc77d720124ba Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?=
 <jmartinezcaamao at gmail.com>
Date: Fri, 7 Nov 2025 12:24:18 +0100
Subject: [PATCH] [SPIRV][SPIRVPrepareGlobals] Map AMD's dynamic LDS 0-element
 globals to arrays with UINT32_MAX elements

In HIP, dynamic LDS globals are represented using 0-element global
arrays in the __shared__ language addressspace.

  extern __shared__ 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 later in AMD's SPIRV runtime
pipeline.
---
 llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp | 27 +++++++++++++++++++
 llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll        | 20 ++++++++++++++
 2 files changed, 47 insertions(+)
 create mode 100644 llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll

diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp
index c44c53129f1e0..42a9577bb2054 100644
--- a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp
@@ -13,6 +13,7 @@
 
 #include "SPIRV.h"
 
+#include "llvm/ADT/STLExtras.h"
 #include "llvm/IR/Module.h"
 
 using namespace llvm;
@@ -43,6 +44,29 @@ bool tryExtendLLVMBitcodeMarker(GlobalVariable &Bitcode) {
   return true;
 }
 
+bool tryExtendDynamicLDSGlobal(GlobalVariable &GV) {
+  constexpr unsigned WorkgroupAS = 3;
+  const bool IsWorkgroupExternal =
+      GV.hasExternalLinkage() && GV.getAddressSpace() == WorkgroupAS;
+  if (!IsWorkgroupExternal)
+    return false;
+
+  const ArrayType *AT = dyn_cast<ArrayType>(GV.getValueType());
+  if (!AT || AT->getNumElements() != 0)
+    return false;
+
+  constexpr auto Magic = std::numeric_limits<uint32_t>::max();
+  ArrayType *NewAT = ArrayType::get(AT->getElementType(), Magic);
+  GlobalVariable *NewGV = new GlobalVariable(
+      *GV.getParent(), NewAT, GV.isConstant(), GV.getLinkage(), nullptr, "",
+      &GV, GV.getThreadLocalMode(), WorkgroupAS, GV.isExternallyInitialized());
+  NewGV->takeName(&GV);
+  GV.replaceAllUsesWith(NewGV);
+  GV.eraseFromParent();
+
+  return true;
+}
+
 bool SPIRVPrepareGlobals::runOnModule(Module &M) {
   const bool IsAMD = M.getTargetTriple().getVendor() == Triple::AMD;
   if (!IsAMD)
@@ -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);
+
   return Changed;
 }
 char SPIRVPrepareGlobals::ID = 0;
diff --git a/llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll b/llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll
new file mode 100644
index 0000000000000..f0acfdfdede9d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll
@@ -0,0 +1,20 @@
+; RUN: llc -verify-machineinstrs -mtriple=spirv64-amd-amdhsa %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -mtriple=spirv64-amd-amdhsa %s -o - -filetype=obj | spirv-val %}
+
+; CHECK: OpName %[[#LDS:]] "lds"
+; CHECK: OpDecorate %[[#LDS]] LinkageAttributes "lds" Import
+; CHECK: %[[#UINT:]] = OpTypeInt 32 0
+; CHECK: %[[#UINT_MAX:]] = OpConstant %[[#UINT]] 4294967295
+; CHECK: %[[#LDS_ARR_TY:]] = OpTypeArray %[[#UINT]] %[[#UINT_MAX]]
+; CHECK: %[[#LDS_ARR_PTR_WG:]] = OpTypePointer Workgroup %[[#LDS_ARR_TY]]
+; CHECK: %[[#LDS]] = OpVariable %[[#LDS_ARR_PTR_WG]] Workgroup
+
+ at lds = external addrspace(3) global [0 x i32]
+
+define spir_kernel void @foo(ptr addrspace(4) %in, ptr addrspace(4) %out) {
+entry:
+  %val = load i32, ptr addrspace(4) %in
+  %add = add i32 %val, 1
+  store i32 %add, ptr addrspace(4) %out
+  ret void
+}