From 9e1d6dde47bd7c2b096a4a1920bdfbbd5e345292 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 9 Jul 2026 16:09:08 +0300 Subject: [PATCH] Lower device globals to kernel arguments (no program-scope globals) rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (#1279). With CHIP_ENABLE_DEVICE_PROGRAM_SCOPE_GLOBALS=OFF, the HipGlobalVariables pass now lowers user __device__/__constant__ globals to implicit trailing kernel pointer arguments: kernels get one pointer arg per referenced global, a __chip_gvararg_ annotation records the mapping for the runtime, init shadow kernels take the storage address as an argument (dispatched by arg count at runtime), and bind shadow kernels are dropped. Globals whose uses cannot be expressed as kernel args (address taken, cross-referencing initializers, uses in non-kernel functions) fall back to program-scope globals. The x86 Intel native CI lane runs the suite with program-scope globals both ON and OFF; unmasks the device-global tests on the rusticl runner. --- .github/workflows/x86-intel-gpu-ci.yml | 10 +- llvm_passes/HipGlobalVariables.cpp | 214 ++++++++++++++++++++++++ scripts/unit_tests.sh | 30 +++- src/CHIPBackend.cc | 31 +++- src/CHIPBackend.hh | 6 + src/CHIPGraph.cc | 15 ++ src/SPIRVFuncInfo.cc | 22 ++- src/SPIRVFuncInfo.hh | 8 + src/backend/Level0/CHIPBackendLevel0.cc | 11 ++ src/backend/OpenCL/CHIPBackendOpenCL.cc | 11 ++ src/common.hh | 8 + src/spirv.cc | 73 ++++++-- tests/known_failures.yaml | 30 +--- 13 files changed, 420 insertions(+), 49 deletions(-) diff --git a/.github/workflows/x86-intel-gpu-ci.yml b/.github/workflows/x86-intel-gpu-ci.yml index 49720fdf7..85826c354 100644 --- a/.github/workflows/x86-intel-gpu-ci.yml +++ b/.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: diff --git a/llvm_passes/HipGlobalVariables.cpp b/llvm_passes/HipGlobalVariables.cpp index 0a58ab7f3..bc368a08d 100644 --- a/llvm_passes/HipGlobalVariables.cpp +++ b/llvm_passes/HipGlobalVariables.cpp @@ -530,6 +530,215 @@ bool emitNonSymbolInitializerKernel(const std::vector 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_` 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_` gets a trailing +// `i8 addrspace(1)*` parameter per distinct G it uses; the load is replaced +// by ptrtoint(param). +// * a `__chip_gvararg_` 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_` shadow kernel is reworked to take the storage +// address as its argument (instead of reading the global), and the +// `__chip_var_bind_` shadow kernel (which only stored into the global) +// is removed. `__chip_var_info_` 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 Loads; + for (User *U : GV->users()) + if (auto *LD = dyn_cast(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 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_` annotation: a constant i8 array holding the +// NUL-separated original global names in trailing-argument order. +static void emitGVarArgAnnotation(Module &M, StringRef KernelName, + ArrayRef 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_` 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(U); + if (!I) + return false; // constant expression use, etc. + Function *F = I->getFunction(); + StringRef FN = F->getName(); + if (isa(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(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 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_` 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 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 Used; + for (auto *GV : GVs) + for (User *U : GV->users()) + if (auto *LD = dyn_cast(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 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_` 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 diff --git a/scripts/unit_tests.sh b/scripts/unit_tests.sh index 7959cba8b..c01fc73fa 100755 --- a/scripts/unit_tests.sh +++ b/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 [--variant=translator|native] [--skip-build] [--build-only] [--num-tries=$num_tries] [--num-threads=$num_threads] [--timeout=$timeout]" + echo "Usage: $0 [--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" } diff --git a/src/CHIPBackend.cc b/src/CHIPBackend.cc index c0bb669fc..44049abad 100755 --- a/src/CHIPBackend.cc +++ b/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) { diff --git a/src/CHIPBackend.hh b/src/CHIPBackend.hh index cf3b7b4b7..ce7564da3 100644 --- a/src/CHIPBackend.hh +++ b/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 HostBuffer_; diff --git a/src/CHIPGraph.cc b/src/CHIPGraph.cc index 103fa4ce8..55fd3c096 100644 --- a/src/CHIPGraph.cc +++ b/src/CHIPGraph.cc @@ -114,6 +114,13 @@ void CHIPGraphNodeMemcpy::execute(chipstar::Queue *Queue) const { } } void CHIPGraphNodeKernel::execute(chipstar::Queue *Queue) const { + // Ensure the kernel module's device variables are allocated before launch. + // The normal hipLaunchKernel path does this, but graph-node execution + // bypasses it — which matters when globals are lowered to kernel arguments + // (their device address must be bound at launch). + if (auto *K = ExecItem_->getKernel()) + if (const void *HPtr = K->getHostPtr()) + Queue->getDevice()->prepareDeviceVariables(HostPtr(HPtr)); Queue->launch(ExecItem_); } @@ -139,6 +146,10 @@ CHIPGraphNodeKernel::CHIPGraphNodeKernel(const hipKernelNodeParams *TheParams) Params_.sharedMemBytes, nullptr); ExecItem_->setKernel(ChipKernel); ExecItem_->setArgs(TheParams->kernelParams); + // setupAllArgs() binds implicit device-global address arguments, so the + // module's device variables must be allocated first. The normal launch path + // does this, but graph-node construction happens before any launch. + Dev->prepareDeviceVariables(HostPtr(Params_.func)); ExecItem_->setupAllArgs(); } @@ -165,6 +176,10 @@ CHIPGraphNodeKernel::CHIPGraphNodeKernel(const void *HostFunction, dim3 GridDim, ExecItem_ = Backend->createExecItem(GridDim, BlockDim, SharedMem, nullptr); ExecItem_->setKernel(ChipKernel); ExecItem_->setArgs(Params_.kernelParams); + // setupAllArgs() binds implicit device-global address arguments, so the + // module's device variables must be allocated first (see the + // hipKernelNodeParams constructor above). + Dev->prepareDeviceVariables(HostPtr(HostFunction)); ExecItem_->setupAllArgs(); } diff --git a/src/SPIRVFuncInfo.cc b/src/SPIRVFuncInfo.cc index 50c89f054..aec35f922 100644 --- a/src/SPIRVFuncInfo.cc +++ b/src/SPIRVFuncInfo.cc @@ -81,6 +81,8 @@ std::string_view SPVFuncInfo::Arg::getKindAsString() const { return "Image"; case SPVTypeKind::Sampler: return "Sampler"; + case SPVTypeKind::DeviceGlobal: + return "DeviceGlobal"; } } @@ -107,6 +109,10 @@ void SPVFuncInfo::visitClientArgsImpl(void **ClientArgList, // <<<>>>-syntax - not in a kernel parameter list. if (ArgTI.isWorkgroupPtr()) continue; + // Implicit device-global address argument (rusticl globals-as-kernel-args + // lowering): provided by the runtime, not visible to the HIP client. + if (ArgKind == SPVTypeKind::DeviceGlobal) + continue; // Map kernel argument types to types as defined in HIP source code. if (ArgKind == SPVTypeKind::Image) @@ -152,19 +158,26 @@ void SPVFuncInfo::visitKernelArgsImpl(void **ClientArgList, if (ArgKind == SPVTypeKind::Sampler) ArgListIndex--; + // DeviceGlobal args are implicit (provided by the runtime, not the client), + // so they don't consume an entry from the client argument list. + bool IsImplicit = + ArgTI.isWorkgroupPtr() || ArgKind == SPVTypeKind::DeviceGlobal; + const void *ArgData = nullptr; - if (ClientArgList && !ArgTI.isWorkgroupPtr()) { + if (ClientArgList && !IsImplicit) { ArgData = ClientArgList[ArgListIndex]; // Clang geerated argument list should not have nullptrs in it. assert(ArgData && "nullptr in the argument list"); } - KernelArg KArg{{{ArgKind, ArgTI.StorageClass, ArgSize}, ArgIndex, ArgData}}; + KernelArg KArg{{{ArgKind, ArgTI.StorageClass, ArgSize, ArgTI.DevGlobalName}, + ArgIndex, ArgData}}; Visitor(KArg); ArgIndex++; - ArgListIndex++; + if (ArgKind != SPVTypeKind::DeviceGlobal) + ArgListIndex++; } } @@ -184,7 +197,8 @@ unsigned SPVFuncInfo::getNumClientArgs() const { unsigned Count = getNumKernelArgs(); for (const auto &ArgTI : ArgTypeInfo_) { auto ArgKind = ArgTI.Kind; - Count -= ArgKind == SPVTypeKind::Sampler || ArgTI.isWorkgroupPtr(); + Count -= ArgKind == SPVTypeKind::Sampler || ArgTI.isWorkgroupPtr() || + ArgKind == SPVTypeKind::DeviceGlobal; } return Count; } diff --git a/src/SPIRVFuncInfo.hh b/src/SPIRVFuncInfo.hh index 778360cbe..cee35a3c2 100644 --- a/src/SPIRVFuncInfo.hh +++ b/src/SPIRVFuncInfo.hh @@ -28,6 +28,7 @@ #include #include #include +#include #include enum class SPVTypeKind : unsigned { @@ -44,6 +45,10 @@ enum class SPVTypeKind : unsigned { // a device buffer. Image, // The type is a image. Sampler, // The type is a sample. + DeviceGlobal, // An implicit trailing pointer argument carrying the device + // address of a __device__/__constant__ global variable (used on + // drivers that can't consume program-scope globals, e.g. + // rusticl). DevGlobalName names the global; no client arg. // Should not appear in kernel parameter lists. Opaque, // The type is an unresolved, special SPIR-V type. @@ -62,6 +67,9 @@ struct SPVArgTypeInfo { SPVTypeKind Kind; SPVStorageClass StorageClass; size_t Size; + /// For Kind==DeviceGlobal: the name of the device global whose address this + /// implicit argument carries. Empty otherwise. + std::string DevGlobalName; bool isWorkgroupPtr() const { return Kind == SPVTypeKind::Pointer && diff --git a/src/backend/Level0/CHIPBackendLevel0.cc b/src/backend/Level0/CHIPBackendLevel0.cc index 3919adc08..bcbec826b 100644 --- a/src/backend/Level0/CHIPBackendLevel0.cc +++ b/src/backend/Level0/CHIPBackendLevel0.cc @@ -3563,6 +3563,17 @@ void CHIPExecItemLevel0::setupAllArgs() { sizeof(void *), &SpillSlot); break; } + case SPVTypeKind::DeviceGlobal: { + // Implicit arg carrying the device address of a __device__/__constant__ + // global (globals-as-kernel-args lowering). Bind it to the global's + // allocated storage. + void *DevPtr = chipstar::getDeviceGlobalArgAddr(Kernel, Arg); + logTrace("setArg {} for device global '{}' -> {}", Arg.Index, + Arg.DevGlobalName, DevPtr); + zeStatus = zeKernelSetArgumentValue(Kernel->get(), Arg.Index, + sizeof(void *), &DevPtr); + break; + } default: CHIPERR_LOG_AND_ABORT( "Internal chipStar error: CHIPExecItemLevel0::setupAllArgs Unknown " diff --git a/src/backend/OpenCL/CHIPBackendOpenCL.cc b/src/backend/OpenCL/CHIPBackendOpenCL.cc index 51c1ebc07..b272764ec 100644 --- a/src/backend/OpenCL/CHIPBackendOpenCL.cc +++ b/src/backend/OpenCL/CHIPBackendOpenCL.cc @@ -2469,6 +2469,17 @@ void CHIPExecItemOpenCL::setupAllArgs() { CHIPERR_CHECK_LOG_AND_THROW_TABLE(clSetKernelArgSVMPointer); break; } + case SPVTypeKind::DeviceGlobal: { + // Implicit arg carrying the device address of a __device__/__constant__ + // global (rusticl globals-as-kernel-args lowering). Bind it to the + // global's allocated storage. + void *DevPtr = chipstar::getDeviceGlobalArgAddr(Kernel, Arg); + logTrace("clSetKernelArgSVMPointer {} for device global '{}' -> {}", + Arg.Index, Arg.DevGlobalName, DevPtr); + Err = ::clSetKernelArgSVMPointer(KernelHandle, Arg.Index, DevPtr); + CHIPERR_CHECK_LOG_AND_THROW_TABLE(clSetKernelArgSVMPointer); + break; + } } }; FuncInfo->visitKernelArgs(getArgs(), ArgVisitor); diff --git a/src/common.hh b/src/common.hh index 3bd9bd337..051486087 100644 --- a/src/common.hh +++ b/src/common.hh @@ -86,6 +86,14 @@ constexpr char ChipNonSymbolResetKernelName[] = "__chip_reset_non_symbols"; /// variables is '' constexpr char ChipSpilledArgsVarPrefix[] = "__chip_spilled_args_"; +/// The prefix for global-scope annotation variables recording which device +/// globals feed a kernel's implicit trailing DeviceGlobal arguments +/// (globals-as-kernel-args lowering, used when program-scope globals are +/// disabled). The variable '' holds the +/// NUL-separated original global names in trailing-argument order. See +/// HipGlobalVariables.cpp for details. +constexpr char ChipGVarArgPrefix[] = "__chip_gvararg_"; + /// The name of a global variable which indicates, when non-zero, if /// the abort() function was called by a kernel. constexpr char ChipDeviceAbortFlagName[] = "__chipspv_abort_called"; diff --git a/src/spirv.cc b/src/spirv.cc index cd2c7708e..7c4e042bc 100644 --- a/src/spirv.cc +++ b/src/spirv.cc @@ -557,6 +557,9 @@ class SPIRVmodule { std::map LinkNames_; std::map>> SpilledArgAnnotations_; + // Kernel-name -> ordered device-global names feeding the trailing implicit + // DeviceGlobal arguments (rusticl globals-as-kernel-args lowering). + std::map> GVarArgAnnotations_; // This flag indicates if the module is known not to have indirect // global buffer accesses (IGBA) in any kernel. This is told by a @@ -629,6 +632,22 @@ class SPIRVmodule { } } + // Mark the trailing implicit DeviceGlobal arguments (rusticl + // globals-as-kernel-args lowering). They are appended after the user + // args, in annotation order. + auto GVarArgsIt = GVarArgAnnotations_.find(KernelName); + if (GVarArgsIt != GVarArgAnnotations_.end()) { + auto &Names = GVarArgsIt->second; + size_t Total = FnInfo->ArgTypeInfo_.size(); + if (Names.size() <= Total) { + size_t Base = Total - Names.size(); + for (size_t J = 0; J < Names.size(); ++J) { + FnInfo->ArgTypeInfo_[Base + J].Kind = SPVTypeKind::DeviceGlobal; + FnInfo->ArgTypeInfo_[Base + J].DevGlobalName = Names[J]; + } + } + } + ModuleInfo.FuncInfoMap.emplace(std::make_pair(i.second, FnInfo)); } KernelInfoMap_.clear(); @@ -652,6 +671,22 @@ class SPIRVmodule { return It != IdToInstMap_.end() ? It->second.get() : nullptr; } + /// Collect the constituent words (the literal in word 3 of each OpConstant + /// element) of an annotation variable's constant-array initializer. + /// 'VarInst' is an OpVariable with an initializer operand (word 4). + std::vector getConstArrayElementWords(const SPIRVinst *VarInst) { + std::vector Words; + auto *Init = getInstruction(VarInst->getWord(4)); + assert(Init && "Annotation variable is missing an initializer."); + auto *Type = TypeMap_[Init->getResultTypeID()]; + assert(Type && dynamic_cast(Type) && + "Could not find type for result ID."); + auto ArrLen = static_cast(Type)->elementCount(); + for (auto EltID : getWordRange(&Init->getWord(3), ArrLen)) + Words.push_back(getInstruction(EltID)->getWord(3)); + return Words; + } + void processKernelParameter(const SPIRVinst &Inst, SPVFuncInfo &FuncInfo) { // Record kernel parameter size for kernel argument setters in the // backends. @@ -782,25 +817,34 @@ class SPIRVmodule { if (startsWith(Name, SpillArgAnnotation) && Inst->size() >= 5) { auto KernelName = Name.substr(SpillArgAnnotation.size()); auto &SpillAnnotation = SpilledArgAnnotations_[KernelName]; - // Get initializer operand (word 4, requires at least 5 words). - auto *Init = getInstruction(Inst->getWord(4)); - assert(Init && "Annotation variable is missing an initializer."); - // Init is known to be OpConstantComposite of char array. - auto *Type = TypeMap_[Init->getResultTypeID()]; - assert(Type && dynamic_cast(Type) && - "Could not type for result ID."); - auto *ArrayType = static_cast(Type); - auto ArrLen = ArrayType->elementCount(); - // Iterate constituents. - for (auto EltID : getWordRange(&Init->getWord(3), ArrLen)) { - auto *ConstInt = getInstruction(EltID); // OpConstant - uint32_t Annotation = ConstInt->getWord(3); + // Annotations are 32-bit words: the lower 16 bits carry the argument + // index of the spilled argument and the upper 16 bits its size. + for (uint32_t Annotation : getConstArrayElementWords(Inst)) { uint16_t ArgIndex = Annotation & 0xffff; uint16_t ArgSize = Annotation >> 16u; SpillAnnotation.push_back(std::make_pair(ArgIndex, ArgSize)); } } + auto GVarArgAnnotation = std::string_view(ChipGVarArgPrefix); + if (startsWith(Name, GVarArgAnnotation) && Inst->size() >= 5) { + auto KernelName = Name.substr(GVarArgAnnotation.size()); + // The initializer is a uchar array holding the NUL-separated + // original global names in trailing-argument order. + std::string Bytes; + for (uint32_t Word : getConstArrayElementWords(Inst)) + Bytes.push_back(static_cast(Word & 0xff)); + // Split on NUL. + auto &Names = GVarArgAnnotations_[KernelName]; + size_t Start = 0; + for (size_t I = 0; I <= Bytes.size(); ++I) + if (I == Bytes.size() || Bytes[I] == '\0') { + if (I > Start) + Names.emplace_back(Bytes.substr(Start, I - Start)); + Start = I + 1; + } + } + if (Name == "__chip_module_has_no_IGBAs" && Inst->size() >= 5) { auto *Init = getInstruction(Inst->getWord(4)); assert(Init && "__chip_module_has_no_IGBAs has invalid initializer ID"); @@ -1027,7 +1071,8 @@ bool postprocessSPIRV(std::vector &Input) { // for mesa/rusticl that does not support them yet. Also, the // variables are essentially dead code for the driver. if (LinkName == "__chip_module_has_no_IGBAs" || - startsWith(LinkName, "__chip_spilled_args_")) + startsWith(LinkName, "__chip_spilled_args_") || + startsWith(LinkName, ChipGVarArgPrefix)) InstructionsToErase.insert(Insn.getWord(1)); } } diff --git a/tests/known_failures.yaml b/tests/known_failures.yaml index 26a19c07c..39d051e37 100644 --- a/tests/known_failures.yaml +++ b/tests/known_failures.yaml @@ -887,25 +887,17 @@ chipstar-rusticl: # AMD Radeon Pro W6400 via rusticl/radeonsi (self-hosted runne cuda-qrng: 'rusticl/radeonsi: CL_INVALID_WORK_GROUP_SIZE / driver gaps on these CUDA-sample ports' cuda-scan: 'rusticl/radeonsi: CL_INVALID_WORK_GROUP_SIZE / driver gaps on these CUDA-sample ports' cuda-vectorAdd: 'rusticl/radeonsi: CL_INVALID_WORK_GROUP_SIZE / driver gaps on these CUDA-sample ports' - hipConstantTestDeviceSymbol: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - hipTestDeviceLink: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - hipTestDeviceSymbol: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - hipTestResetStaticVar: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - hipTestSymbolInit: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - hipTestSymbolReset: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - hipTestVariableTemplateSymbols: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' + hipTestResetStaticVar: 'rusticl/radeonsi: function-local static device variable is not convertible to a kernel argument (issue #1279 kernel-arg lowering), so it falls back to a program-scope CrossWorkgroup global, which rusticl/Mesa cannot consume (driver panic)' + hipTestSymbolInit: 'rusticl/radeonsi: cross-referenced/indirect device globals are not convertible to kernel arguments (issue #1279 kernel-arg lowering), so they fall back to program-scope CrossWorkgroup globals, which rusticl/Mesa cannot consume (driver panic)' PrintfDynamic: 'unsupported on rusticl/radeonsi (driver gap; see rusticl AMD GPU CI)' TestAlignAttrRuntime: 'unsupported on rusticl/radeonsi (driver gap; see rusticl AMD GPU CI)' - TestAtomics: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' + TestAtomics: 'rusticl/ACO clLinkProgram failure (-17): unsupported int64/system-scope atomics capabilities linked in the same module (same root cause as Unit_deviceFunctions_CompileTest_atomicAdd_*; not a device-global issue)' TestBufferDevAddr: 'rusticl lacks cl_ext_buffer_device_address' - TestConstantMemory: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' TestDefaultStreamImplicitSync: 'unsupported on rusticl/radeonsi (driver gap; see rusticl AMD GPU CI)' - TestGlobalVarInit: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - TestLargeGlobalVar: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' + TestLargeGlobalVar: 'rusticl/radeonsi: large device-global test times out (>180s) on the W6400' TestRDCWithMultipleHipccCmds: 'rusticl lacks relocatable-device-code / cross-module device global linking' TestRDCWithSingleHipccCmd: 'rusticl lacks relocatable-device-code / cross-module device global linking' TestSeparateCompilation: 'rusticl lacks relocatable-device-code / cross-module device global linking' - TestTemplatedConstantMemcpy: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' TestWholeProgramCompilation: 'rusticl lacks relocatable-device-code / cross-module device global linking' Unit_deviceFunctions_CompileTest_atomicAdd_float: 'rusticl/ACO clLinkProgram failure (-17): unsupported int64/system-scope atomics (and fp64) capabilities linked in the same module' Unit_deviceFunctions_CompileTest_atomicAdd_int: 'rusticl/ACO clLinkProgram failure (-17): unsupported int64/system-scope atomics (and fp64) capabilities linked in the same module' @@ -916,14 +908,9 @@ chipstar-rusticl: # AMD Radeon Pro W6400 via rusticl/radeonsi (self-hosted runne Unit_deviceFunctions_CompileTest_atomicAdd_unsigned_long_long: 'rusticl/ACO clLinkProgram failure (-17): unsupported int64/system-scope atomics (and fp64) capabilities linked in the same module' Unit_deviceFunctions_CompileTest_atomicAdd_usigned_int: 'rusticl/ACO clLinkProgram failure (-17): unsupported int64/system-scope atomics (and fp64) capabilities linked in the same module' Unit_deviceFunctions_CompileTest_nanf_float: 'rusticl/ACO clLinkProgram failure (-17): unsupported int64/system-scope atomics (and fp64) capabilities linked in the same module' - Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalConstMemory: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemory: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryWithKernel: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - Unit_hipGraphAddMemcpyNodeToSymbol_GlobalConstMemory: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - Unit_hipGraphAddMemcpyNodeToSymbol_GlobalMemory: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - Unit_hipGraphAddMemcpyNodeToSymbol_MemcpyToSymbolNodeWithKernel: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Functional: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - Unit_hipGraphMemcpyNodeSetParamsToSymbol_Functional: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' + Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalConstMemory: 'rusticl/radeonsi: flaky only under the full -j4 CI suite -- the graph memcpy-to-symbol -> memcpy-from-symbol round-trip intermittently reads stale data on the W6400 under concurrent load. The kernel-argument lowering itself is correct: passes reliably in isolation on rusticl and on Intel iGPU / Level Zero. Suspected rusticl driver ordering gap under load (issue #1279 follow-up)' + Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemory: 'rusticl/radeonsi: flaky only under the full -j4 CI suite -- the graph memcpy-to-symbol -> memcpy-from-symbol round-trip intermittently reads stale data on the W6400 under concurrent load. The kernel-argument lowering itself is correct: passes reliably in isolation on rusticl and on Intel iGPU / Level Zero. Suspected rusticl driver ordering gap under load (issue #1279 follow-up)' + Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryWithKernel: 'rusticl/radeonsi: flaky only under the full -j4 CI suite -- the graph memcpy-to-symbol -> memcpy-from-symbol round-trip intermittently reads stale data on the W6400 under concurrent load. The kernel-argument lowering itself is correct: passes reliably in isolation on rusticl and on Intel iGPU / Level Zero. Suspected rusticl driver ordering gap under load (issue #1279 follow-up)' Unit_hipHostMalloc_CoherentAccess: 'unsupported on rusticl/radeonsi (driver gap; see rusticl AMD GPU CI)' Unit_hipMallocManaged_AccessMultiStream: 'rusticl is coarse-grain-SVM only; no managed/unified memory (hipMallocManaged/prefetch/advise)' Unit_hipMallocManaged_FlgParam: 'rusticl is coarse-grain-SVM only; no managed/unified memory (hipMallocManaged/prefetch/advise)' @@ -932,13 +919,10 @@ chipstar-rusticl: # AMD Radeon Pro W6400 via rusticl/radeonsi (self-hosted runne Unit_hipMallocManaged_TwoPointers - float: 'rusticl is coarse-grain-SVM only; no managed/unified memory (hipMallocManaged/prefetch/advise)' Unit_hipMallocManaged_TwoPointers - int: 'rusticl is coarse-grain-SVM only; no managed/unified memory (hipMallocManaged/prefetch/advise)' Unit_hipMemAdvise_TstAccessedByFlg4: 'rusticl is coarse-grain-SVM only; no managed/unified memory (hipMallocManaged/prefetch/advise)' - Unit_hipMemcpyFromToSymbol_Negative: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' - Unit_hipMemcpyFromToSymbol_Negative_MemoryTest2: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (same root cause as Unit_hipMemcpyFromToSymbol_Negative); negative test aborts with "Destination is nullptr". Surfaced by narrowing the over-broad _MemoryTest exclusion.' Unit_hipMemcpy2D_Positive_Basic_MemoryTest1: 'rusticl/radeonsi advertises maxThreadsPerBlock=1024 but the fill-kernel work-group limit is lower, so the buffer-fill launch is rejected with CL_INVALID_WORK_GROUP_SIZE even after clamping to the reported device max. Passes on devices that report their real per-block limit (Intel iGPU, PVC).' Unit_hipMemcpy_Positive_Basic_MemoryTest2: 'rusticl/radeonsi advertises maxThreadsPerBlock=1024 but the fill-kernel work-group limit is lower, so the buffer-fill launch is rejected with CL_INVALID_WORK_GROUP_SIZE even after clamping to the reported device max. Passes on devices that report their real per-block limit (Intel iGPU, PVC).' Unit_hipMemcpyDtoH_Positive_Basic_MemoryTest2: 'rusticl/radeonsi advertises maxThreadsPerBlock=1024 but the fill-kernel work-group limit is lower, so the buffer-fill launch is rejected with CL_INVALID_WORK_GROUP_SIZE even after clamping to the reported device max. Passes on devices that report their real per-block limit (Intel iGPU, PVC).' Unit_hipMemcpyDtoHAsync_Positive_Basic_MemoryTest2: 'rusticl/radeonsi advertises maxThreadsPerBlock=1024 but the fill-kernel work-group limit is lower, so the buffer-fill launch is rejected with CL_INVALID_WORK_GROUP_SIZE even after clamping to the reported device max. Passes on devices that report their real per-block limit (Intel iGPU, PVC).' - Unit_hipMemcpyToFromSymbol_SyncAndAsync: 'rusticl/radeonsi cannot consume program-scope CrossWorkgroup globals (device/__constant__ globals, symbols, static vars): "Initializer for CrossWorkgroup variable not yet supported in Mesa"' Unit_hipMemFaultStackAllocation_Check: 'unsupported on rusticl/radeonsi (driver gap; see rusticl AMD GPU CI)' Unit_hipMemPrefetchAsync_NonPageSz: 'rusticl is coarse-grain-SVM only; no managed/unified memory (hipMallocManaged/prefetch/advise)' 'x4\d\d\dc\ds\db0n0': # aurora