Lower device globals to kernel arguments (no program-scope globals)

.github/workflows/x86-intel-gpu-ci.yml

@@ -68,7 +68,10 @@ jobs:
           fi
           echo "OK: refused as expected"
 
-  # Stage 1: Build and test chipStar with latest LLVM native
+  # Stage 1: Build and test chipStar with latest LLVM native.
+  # Runs the suite twice: once with program-scope device globals ON (default)
+  # and once OFF. The OFF build lowers user __device__/__constant__ globals to
+  # kernel arguments (issue #1279); a single native run gates both code paths.
   chipstar-llvm22-native-release:
     needs: refuse-chipfft-chipblas-on-level-zero
     runs-on: [self-hosted, Linux, X64]
@@ -80,9 +83,12 @@ jobs:
           ref: ${{ github.event.pull_request.head.sha || github.sha }}
           fetch-depth: 0
           submodules: 'recursive'
-      - name: Run unit test checking script
+      - name: Run unit test checking script (program-scope globals ON)
         run: ./scripts/unit_tests.sh release llvm-22 --variant=native
         shell: bash
+      - name: Run unit test checking script (program-scope globals OFF / kernel-arg path)
+        run: ./scripts/unit_tests.sh release llvm-22 --variant=native --no-psg
+        shell: bash
 
   # Stage 2: Build chipStar + all libraries, run library tests
   build-and-test-libraries:

llvm_passes/HipGlobalVariables.cpp

@@ -530,6 +530,215 @@ bool emitNonSymbolInitializerKernel(const std::vector<GlobalVariable *> GVs,
   return true;
 }
 
+#ifndef CHIP_ENABLE_DEVICE_PROGRAM_SCOPE_GLOBALS
+// ===========================================================================
+// rusticl/radeonsi path: lower device globals to implicit kernel arguments.
+//
+// Some OpenCL drivers (rusticl/radeonsi on Mesa/ACO) cannot consume
+// program-scope CrossWorkgroup globals. The default lowering above leaves an
+// i64 `__chip_var_<name>` address-holder global which trips them. Here we run a
+// post-transform that removes those globals and instead passes each global's
+// device address as an implicit trailing kernel pointer argument:
+//
+//   * every kernel that loads `__chip_var_<G>` gets a trailing
+//     `i8 addrspace(1)*` parameter per distinct G it uses; the load is replaced
+//     by ptrtoint(param).
+//   * a `__chip_gvararg_<kernel>` annotation (NUL-separated original global
+//     names, in parameter order) tells the runtime which global feeds each
+//     trailing arg. spirv.cc strips this annotation before the driver sees it.
+//   * the `__chip_var_init_<G>` shadow kernel is reworked to take the storage
+//     address as its argument (instead of reading the global), and the
+//     `__chip_var_bind_<G>` shadow kernel (which only stored into the global)
+//     is removed. `__chip_var_info_<G>` is left intact.
+//
+// The appended args must remain the TRAILING kernel args: the runtime marks
+// the last N args as DeviceGlobal by position. HipKernelArgSpiller — the only
+// later pass modifying kernel parameter lists — replaces args in place
+// without changing their count or positions.
+// ===========================================================================
+
+// Replace, inside function F, all LoadInst that load `GV` with ptrtoint(NewPtr).
+static void replaceGlobalLoadsWith(Function &F, GlobalVariable *GV,
+                                   Value *NewPtr) {
+  IRBuilder<> B(F.getContext());
+  SmallVector<LoadInst *, 4> Loads;
+  for (User *U : GV->users())
+    if (auto *LD = dyn_cast<LoadInst>(U))
+      if (LD->getFunction() == &F)
+        Loads.push_back(LD);
+  for (auto *LD : Loads) {
+    B.SetInsertPoint(LD);
+    Value *AsInt = B.CreatePtrToInt(NewPtr, LD->getType());
+    LD->replaceAllUsesWith(AsInt);
+    LD->eraseFromParent();
+  }
+}
+
+// Clone F into a new function with `NumExtra` extra trailing i8 addrspace(1)*
+// parameters, moving the body over. Returns the new function (old one erased).
+// The caller is responsible for wiring the extra args.
+static Function *appendPtrArgs(Function *F, unsigned NumExtra) {
+  LLVMContext &C = F->getContext();
+  auto *PtrTy = PointerType::get(C, SpirvCrossWorkGroupAS);
+  SmallVector<Type *> ArgTys;
+  for (auto &A : F->args())
+    ArgTys.push_back(A.getType());
+  for (unsigned i = 0; i < NumExtra; ++i)
+    ArgTys.push_back(PtrTy);
+  auto *NewFT = FunctionType::get(F->getReturnType(), ArgTys, F->isVarArg());
+  auto *NF = Function::Create(NewFT, F->getLinkage(), F->getAddressSpace(), "",
+                              F->getParent());
+  NF->setCallingConv(F->getCallingConv());
+  NF->setVisibility(F->getVisibility());
+  NF->setAttributes(F->getAttributes()); // function + existing-arg attributes
+  NF->copyMetadata(F, 0); // function-level metadata (!dbg, !reqd_* etc.)
+  NF->takeName(F);
+  // Move the body across and rewire the original arguments.
+  NF->splice(NF->begin(), F);
+  for (unsigned i = 0; i < F->arg_size(); ++i) {
+    F->getArg(i)->replaceAllUsesWith(NF->getArg(i));
+    NF->getArg(i)->takeName(F->getArg(i));
+  }
+  F->eraseFromParent();
+  return NF;
+}
+
+// Emit the `__chip_gvararg_<kernel>` annotation: a constant i8 array holding the
+// NUL-separated original global names in trailing-argument order.
+static void emitGVarArgAnnotation(Module &M, StringRef KernelName,
+                                  ArrayRef<std::string> GlobalNames) {
+  std::string Buf;
+  for (auto &N : GlobalNames) {
+    Buf += N;
+    Buf += '\0';
+  }
+  auto *Init = ConstantDataArray::getString(M.getContext(), Buf,
+                                             /*AddNull=*/false);
+  auto Name = (Twine(ChipGVarArgPrefix) + KernelName).str();
+  new GlobalVariable(M, Init->getType(), /*isConstant=*/true,
+                     GlobalValue::ExternalLinkage, Init, Name, nullptr,
+                     GlobalValue::NotThreadLocal, SpirvCrossWorkGroupAS);
+}
+
+// Original device-global name from a `__chip_var_<name>` address holder.
+static std::string originalNameOf(const GlobalVariable *NewGV) {
+  StringRef N = NewGV->getName();
+  N.consume_front(ChipVarPrefix);
+  return N.str();
+}
+
+// Returns true if F is a chipStar-generated shadow kernel (info/bind/init or
+// the non-symbol reset kernel).
+static bool isShadowKernel(const Function &F) {
+  StringRef N = F.getName();
+  return N.starts_with(ChipVarInfoPrefix) || N.starts_with(ChipVarBindPrefix) ||
+         N.starts_with(ChipVarInitPrefix) || N == ChipNonSymbolResetKernelName;
+}
+
+// A lowered address-holder global can be safely converted to kernel arguments
+// only if every use is either a load inside a user (SPIR_KERNEL) kernel, a load
+// inside its OWN init shadow kernel, or the store inside its OWN bind shadow
+// kernel. Anything else — the address taken (ptrtoint/GEP/constexpr), a load in
+// a non-kernel device function, or a cross-reference from another global's init
+// kernel (e.g. mutually-referencing pointers `Foo=&Bar; Bar=&Foo`) — cannot be
+// expressed as a per-global trailing argument, so we leave such globals as
+// program-scope (which still works on drivers that support them).
+static bool isConvertibleGlobal(GlobalVariable *GV) {
+  std::string OrigName = originalNameOf(GV);
+  std::string OwnInit = (Twine(ChipVarInitPrefix) + OrigName).str();
+  std::string OwnBind = (Twine(ChipVarBindPrefix) + OrigName).str();
+  for (User *U : GV->users()) {
+    auto *I = dyn_cast<Instruction>(U);
+    if (!I)
+      return false; // constant expression use, etc.
+    Function *F = I->getFunction();
+    StringRef FN = F->getName();
+    if (isa<LoadInst>(I)) {
+      if (FN == OwnInit)
+        continue;
+      if (F->getCallingConv() == CallingConv::SPIR_KERNEL && !isShadowKernel(*F))
+        continue; // load in a user kernel
+      return false; // cross-ref init, or load in a non-kernel function
+    }
+    if (isa<StoreInst>(I)) {
+      if (FN == OwnBind)
+        continue; // own bind kernel's store
+      return false;
+    }
+    return false; // any other use kind
+  }
+  return true;
+}
+
+static bool lowerGlobalsToKernelArgs(Module &M, GVarMapT &GVarMap) {
+  if (GVarMap.empty())
+    return false;
+
+  // The lowered address-holder globals we can safely convert to kernel args.
+  // Non-convertible ones are left as program-scope globals (driver fallback).
+  SmallVector<GlobalVariable *, 8> GVs;
+  for (auto &Kv : GVarMap)
+    if (isConvertibleGlobal(Kv.second))
+      GVs.push_back(Kv.second);
+  if (GVs.empty())
+    return false;
+
+  // 1) Rework each `__chip_var_init_<G>` shadow kernel to take the storage
+  //    address as a trailing pointer arg instead of loading the global.
+  for (auto *GV : GVs) {
+    std::string OrigName = originalNameOf(GV);
+    auto *InitF = M.getFunction((Twine(ChipVarInitPrefix) + OrigName).str());
+    if (!InitF || InitF->isDeclaration())
+      continue;
+    Function *NInit = appendPtrArgs(InitF, 1);
+    replaceGlobalLoadsWith(*NInit, GV, NInit->getArg(NInit->arg_size() - 1));
+  }
+
+  // 2) Transform user kernels: append a pointer arg per global they load.
+  SmallVector<Function *, 16> UserKernels;
+  for (auto &F : M)
+    if (F.getCallingConv() == CallingConv::SPIR_KERNEL && !F.isDeclaration() &&
+        !isShadowKernel(F))
+      UserKernels.push_back(&F);
+
+  for (Function *F : UserKernels) {
+    // Collect, in deterministic order, the globals this kernel loads (GVs is
+    // duplicate-free, so Used is too).
+    SmallVector<GlobalVariable *, 4> Used;
+    for (auto *GV : GVs)
+      for (User *U : GV->users())
+        if (auto *LD = dyn_cast<LoadInst>(U))
+          if (LD->getFunction() == F) {
+            Used.push_back(GV);
+            break;
+          }
+    if (Used.empty())
+      continue;
+
+    Function *NF = appendPtrArgs(F, Used.size());
+    unsigned Base = NF->arg_size() - Used.size();
+    std::vector<std::string> Names;
+    for (unsigned j = 0; j < Used.size(); ++j) {
+      replaceGlobalLoadsWith(*NF, Used[j], NF->getArg(Base + j));
+      Names.push_back(originalNameOf(Used[j]));
+    }
+    emitGVarArgAnnotation(M, NF->getName(), Names);
+  }
+
+  // 3) Remove the `__chip_var_bind_<G>` shadow kernels (now pointless) and the
+  //    address-holder globals themselves. Erasing the bind kernel first drops
+  //    the only remaining use of the global.
+  for (auto *GV : GVs) {
+    std::string OrigName = originalNameOf(GV);
+    if (auto *BindF = M.getFunction((Twine(ChipVarBindPrefix) + OrigName).str()))
+      BindF->eraseFromParent();
+    GV->replaceAllUsesWith(PoisonValue::get(GV->getType()));
+    GV->eraseFromParent();
+  }
+  return true;
+}
+#endif // !CHIP_ENABLE_DEVICE_PROGRAM_SCOPE_GLOBALS
+
 static bool lowerGlobalVariables(Module &M) {
   bool Changed = false;
 
@@ -552,6 +761,11 @@ static bool lowerGlobalVariables(Module &M) {
     replaceGlobalVariableUses(GVarMap);
     eraseMappedGlobalVariables(GVarMap);
     Changed |= true;
+#ifndef CHIP_ENABLE_DEVICE_PROGRAM_SCOPE_GLOBALS
+    // rusticl path: convert the lowered address-holder globals to implicit
+    // kernel arguments (the globals themselves are removed here).
+    lowerGlobalsToKernelArgs(M, GVarMap);
+#endif
   }
 
   // Lower global device variables which are not accessible by the host but

scripts/unit_tests.sh

@@ -21,12 +21,15 @@ num_tries=1
 # deafult timeout is 30 minutes
 timeout=1800
 build_only=false
+# When true, build with program-scope device globals disabled so user
+# __device__/__constant__ globals are lowered to kernel arguments instead.
+psg_off=false
 
 rm -rf ~/.cache/chipStar
 
 # Check if at least one argument is provided
 if [ "$#" -lt 2 ]; then
-  echo "Usage: $0 <debug|release> <llvm-20|llvm-21|llvm-22> [--variant=translator|native] [--skip-build] [--build-only] [--num-tries=$num_tries] [--num-threads=$num_threads] [--timeout=$timeout]"
+  echo "Usage: $0 <debug|release> <llvm-20|llvm-21|llvm-22> [--variant=translator|native] [--skip-build] [--build-only] [--no-psg] [--num-tries=$num_tries] [--num-threads=$num_threads] [--timeout=$timeout]"
   exit 1
 fi
 
@@ -89,6 +92,10 @@ do
       build_only=true
       shift
       ;;
+    --no-psg)
+      psg_off=true
+      shift
+      ;;
     --timeout=*)
       timeout="${arg#*=}"
       shift
@@ -142,6 +149,7 @@ echo "num_tries   = ${num_tries}"
 echo "num_threads = ${num_threads}"
 echo "skip_build  = ${skip_build}"
 echo "build_only  = ${build_only}"
+echo "psg_off     = ${psg_off}"
 echo "timeout     = ${timeout}"
 
 # source /opt/intel/oneapi/setvars.sh intel64 &> /dev/null
@@ -199,6 +207,13 @@ else
   else
     CHIP_OPTIONS="-DCHIP_BUILD_SAMPLES=ON -DCHIP_BUILD_TESTS=ON"
   fi
+  # Exercise the kernel-argument lowering path for device globals: with
+  # program-scope globals OFF, user __device__/__constant__ globals become
+  # implicit kernel pointer arguments instead of program-scope CrossWorkgroup
+  # globals (issue #1279).
+  if [ "$psg_off" = true ]; then
+    CHIP_OPTIONS="${CHIP_OPTIONS} -DCHIP_ENABLE_DEVICE_PROGRAM_SCOPE_GLOBALS=OFF"
+  fi
   # Build the project
   echo "Building project..."
   rm -rf HIPCC
@@ -229,12 +244,23 @@ fi
 
 module unload opencl/dgpu
 
+# Tests that busy-wait on clock64() in device code (e.g. Unit_hipHostMalloc_CoherentAccess
+# spins `do { cur = clock64()/clkRate - start; } while (cur < wait_sec);`). With program-scope
+# device globals OFF the clock counter (__chip_clk_counter) is omitted and clock64() returns 0,
+# so the loop never terminates: the kernel hangs and the GPU watchdog aborts the queue with
+# CL_OUT_OF_RESOURCES. These tests use a feature that is intentionally unavailable in the
+# --no-psg build, so skip them there only; they run normally with program-scope globals ON.
+PSG_OFF_EXCLUDE=""
+if [ "$psg_off" = true ]; then
+  PSG_OFF_EXCLUDE='--regex-exclude=Unit_hipHostMalloc_CoherentAccess$'
+fi
+
 # Function to run tests
 run_tests() {
     local device=$1
     local backend=$2
     echo "begin ${device}_${backend}_failed_tests"
-    ../scripts/check.py ./ $device $backend --num-threads=${num_threads} --timeout=$timeout --num-tries=$num_tries | tee ${device}_${backend}_make_check_result.txt
+    ../scripts/check.py ./ $device $backend ${PSG_OFF_EXCLUDE} --num-threads=${num_threads} --timeout=$timeout --num-tries=$num_tries | tee ${device}_${backend}_make_check_result.txt
     echo "end ${device}_${backend}_failed_tests"
 }
 

src/CHIPBackend.cc

@@ -78,9 +78,14 @@ static void queueVariableBindShadowKernel(chipstar::Queue *Q,
   assert(M && Var);
   auto *DevPtr = Var->getDevAddr();
   assert(DevPtr && "Space has not be allocated for a variable.");
-  auto *K = M->getKernelByName(std::string(ChipVarBindPrefix) +
-                               std::string(Var->getName()));
-  assert(K && "chipstar::Module is missing a shadow kernel?");
+  // Use the non-throwing findKernel(): when device globals are lowered to
+  // kernel arguments (rusticl path), the bind shadow kernels are removed
+  // because there is no program-scope global to bind — the address travels as
+  // an implicit kernel argument instead.
+  auto *K = M->findKernel(std::string(ChipVarBindPrefix) +
+                          std::string(Var->getName()));
+  if (!K)
+    return;
   void *Args[] = {&DevPtr};
   queueKernel(Q, K, Args);
 }
@@ -92,7 +97,25 @@ static void queueVariableInitShadowKernel(chipstar::Queue *Q,
   auto *K = M->getKernelByName(std::string(ChipVarInitPrefix) +
                                std::string(Var->getName()));
   assert(K && "chipstar::Module is missing a shadow kernel?");
-  queueKernel(Q, K);
+  if (K->getFuncInfo()->getNumKernelArgs() == 1) {
+    // Globals-as-kernel-args lowering: the init kernel takes the storage
+    // address as its argument instead of reading a program-scope global.
+    auto *DevPtr = Var->getDevAddr();
+    void *Args[] = {&DevPtr};
+    queueKernel(Q, K, Args);
+  } else
+    queueKernel(Q, K);
+}
+
+void *chipstar::getDeviceGlobalArgAddr(chipstar::Kernel *Kernel,
+                                       const SPVFuncInfo::KernelArg &Arg) {
+  auto *Var = Kernel->getModule()->getGlobalVar(Arg.DevGlobalName.c_str());
+  if (!Var || !Var->getDevAddr())
+    CHIPERR_LOG_AND_THROW(
+        "DeviceGlobal kernel arg references an unallocated global: " +
+            Arg.DevGlobalName,
+        hipErrorLaunchFailure);
+  return Var->getDevAddr();
 }
 
 static void initDeviceHeap(chipstar::Queue *Q, chipstar::Module *M) {

src/CHIPBackend.hh

@@ -1163,6 +1163,12 @@ public:
   virtual const chipstar::Module *getModule() const = 0;
 };
 
+/// Resolve the storage address of the device global backing an implicit
+/// DeviceGlobal kernel argument (globals-as-kernel-args lowering). Throws
+/// hipErrorLaunchFailure if the global's storage is not allocated.
+void *getDeviceGlobalArgAddr(chipstar::Kernel *Kernel,
+                             const SPVFuncInfo::KernelArg &Arg);
+
 class ArgSpillBuffer {
   chipstar::Context *Ctx_; ///< A context to allocate device space from.
   std::unique_ptr<char[]> HostBuffer_;