From 87af0be6c38fed61c315980eb165e17f01862ad0 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 17 Sep 2026 05:44:32 +0300 Subject: [PATCH 1/2] tests: add reproducer for HipDynMem clone order (#1661) Running hip-post-link-passes twice on the same module can give different output when a dynamic shared memory global (extern __shared__ T x[]) has more than 16 direct users: HipDynMemExternReplacePass collects them into a SmallPtrSet, which past its inline capacity iterates in pointer hash order, and it appends one clone per user to the module in that order. Function objects land at different heap addresses in every process under ASLR, so the order of the lowered kernels, and the SPIR-V emitted from them, changes from run to run for an unchanged source. Caches keyed on the emitted IL (the OpenCL module cache, PoCL's kernel cache) then miss. The test generates 32 spir_kernel functions that each store to one external addrspace(3) [0 x i32] global, runs opt with the chipStar plugin and -passes=hip-post-link-passes eight times, and requires every textual output to match the first. The module is generated in the script, so there is no fixture file. 32 users puts the set well past its inline size; with 16 or fewer it iterates in insertion order and cannot show the bug. Several runs are needed because the order only changes between processes. Measured on main (7db0385783bd, LLVM 22, x86_64 Linux, ASLR on): the test failed 200 of 200 invocations. Run 2 already differed from run 1 in 192 of them and run 3 in the other 8, so two runs agree about 4% of the time and eight agreeing by chance is not a realistic outcome. On the reduction sample that motivated the issue, 20 runs gave 11 distinct modules. --- tests/compiler/CMakeLists.txt | 2 ++ tests/compiler/TestFix1661DynMemOrder.bash | 31 ++++++++++++++++++++++ 2 files changed, 33 insertions(+) create mode 100644 tests/compiler/TestFix1661DynMemOrder.bash diff --git a/tests/compiler/CMakeLists.txt b/tests/compiler/CMakeLists.txt index 85c355d42..2d54704ee 100644 --- a/tests/compiler/CMakeLists.txt +++ b/tests/compiler/CMakeLists.txt @@ -202,6 +202,8 @@ add_hipcc_test(TestDeviceVirtualFunctions.hip HIPCC_OPTIONS -c) add_shell_test(TestDeadVTableVirtualBase.bash) # Program-scope variable init kernels must use their work items (#582). add_shell_test(TestFix582VarInitWorkItems.bash) +# Lowering a module with many shared memory users must be reproducible (#1661). +add_shell_test(TestFix1661DynMemOrder.bash) add_hipcc_test(TestHostSideHIPVectors.hip HIPCC_OPTIONS -fsyntax-only) add_hipcc_test(TestAlignAttr.hip HIPCC_OPTIONS -fsyntax-only) # Check CHIP_FAST_MATH is set for -ffast-math and preprocessor guards diff --git a/tests/compiler/TestFix1661DynMemOrder.bash b/tests/compiler/TestFix1661DynMemOrder.bash new file mode 100644 index 000000000..2d304836f --- /dev/null +++ b/tests/compiler/TestFix1661DynMemOrder.bash @@ -0,0 +1,31 @@ +#!/bin/bash +# Regression test for CHIP-SPV/chipStar#1661: lowering a module must not depend +# on allocation addresses, so repeated post-link pipeline runs agree. +set -eu + +OPT="@LLVM_TOOLS_BINARY_DIR@/opt" +PLUGIN="@CMAKE_BINARY_DIR@/lib/libLLVMHipSpvPasses.so" +OUT="@CMAKE_CURRENT_BINARY_DIR@/@TEST_NAME@.d" + +rm -rf "${OUT}"; mkdir -p "${OUT}"; cd "${OUT}" + +# More than 16 users: a small pointer set iterates in insertion order anyway. +{ + echo '@__smem = external addrspace(3) global [0 x i32]' + for i in $(seq 32); do + printf 'define spir_kernel void @k%d(i32 %%v) {\n' "${i}" + printf ' store i32 %%v, ptr addrspace(3) @__smem\n ret void\n}\n' + done +} > in.ll + +# Heap addresses only change between processes, so compare several runs. +for i in $(seq 8); do + "${OPT}" -load-pass-plugin "${PLUGIN}" -passes=hip-post-link-passes \ + in.ll -S -o "run${i}.ll" + if ! cmp -s run1.ll "run${i}.ll"; then + echo "FAIL: run ${i} lowered the same module differently from run 1" + diff run1.ll "run${i}.ll" | head -10 + exit 1 + fi +done +echo "PASSED" From 2fe5a4d8e392333ac18e517d59a8a6a0edabc9ec Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 17 Sep 2026 05:45:59 +0300 Subject: [PATCH 2/2] HipDynMem: collect direct users in insertion order (#1661) HipDynMemExternReplacePass gathered the functions that use a dynamic shared memory global into FSet, a SmallPtrSet. That set iterates in insertion order only while it holds 16 or fewer elements; past that it switches to a hash table keyed on the pointer value, so iteration follows heap addresses. The loop over it decides both the order in which indirect users are discovered and the order in which each direct user is cloned with the extra argument, and every clone is appended to the end of the module by Function::Create. With ASLR, two opt runs on the same bitcode therefore produced the same functions in a different order, and different SPIR-V, whenever a global had more than 16 direct users. The cuda-reduction sample has 88 and 44. Collect the direct users into OrderedFSet (llvm::SetVector), the type the pass already uses for indirect users. Iteration is then in use-list order, which depends only on the input module. Membership and the set of clones are unchanged; only their order in the module is fixed. FSet and the SmallPtrSet include had no other users and go with it. The rest of the pass does not depend on addresses: the global variables come from the ValueSymbolTable, a StringMap hashed on names; users are walked through use-lists; the ValueToValueMap given to CloneFunctionInto is only looked up, never iterated for output. Measured with LLVM 22 on x86_64 Linux: TestFix1661DynMemOrder passed 200 of 200 invocations (1600 opt runs), and 30 runs of hip-post-link-passes on reduction_kernel-hip-spirv64-generic-link.bc gave one module, with the same set of defined functions as before the change. Before it, the test failed 200 of 200 and 20 runs of the sample gave 11 distinct modules. --- llvm_passes/HipDynMem.cpp | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/llvm_passes/HipDynMem.cpp b/llvm_passes/HipDynMem.cpp index b4c0258f7..2fd4c9442 100644 --- a/llvm_passes/HipDynMem.cpp +++ b/llvm_passes/HipDynMem.cpp @@ -19,7 +19,6 @@ #include "HipDynMem.h" -#include "llvm/ADT/SmallPtrSet.h" #include "llvm/ADT/SetVector.h" #include "llvm/ADT/Statistic.h" #include "llvm/IR/Function.h" @@ -44,7 +43,6 @@ using namespace llvm; #define SPIR_LOCAL_AS 3 #define GENERIC_AS 4 -typedef llvm::SmallPtrSet FSet; typedef llvm::SetVector OrderedFSet; typedef llvm::SmallVector GVarVec; @@ -91,7 +89,7 @@ class HipDynMemExternReplacePass : public ModulePass { } } - static void recursivelyFindDirectUsers(Value *V, FSet &FS) { + static void recursivelyFindDirectUsers(Value *V, OrderedFSet &FS) { for (auto U : V->users()) { Instruction *Inst = dyn_cast(U); if (Inst) { @@ -440,7 +438,7 @@ CloneFunctionInto(NewF, F, VV, CloneFunctionChangeType::GlobalChanges, RI); for (GlobalVariable *GV : GVars) { - FSet DirectUserSet; + OrderedFSet DirectUserSet; // first, find functions that directly use the GVar. However, these may be // called from other functions, so we need to append the