Conversation
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<Function *, 16>, 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 (7db0385, 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.
HipDynMemExternReplacePass gathered the functions that use a dynamic shared memory global into FSet, a SmallPtrSet<Function *, 16>. 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<Function *>), 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.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
HipDynMemExternReplacePass collected the direct users of a dynamic shared memory global in a SmallPtrSet, which iterates in pointer hash order once it holds more than 16 elements, and appended a clone of each user to the module in that order. Compiling the same source twice could then emit the kernels in a different order, so caches keyed on the SPIR-V missed. The direct users now go into the SetVector the pass already uses for indirect users.
The new compile-only test runs the post-link pipeline several times on a module with 32 users and requires identical output. It relies on heap addresses changing between processes, so it cannot fail where ASLR is disabled.
Fixes #1661
This PR was generated using AI