HipGlobalVariables: expand lowered globals nested in constant aggregates - #1651
Merged
Merged
Conversation
…ce global's address
A device global whose struct initializer holds the address of another
lowered device global reads back a wrong pointer on the device, because
HipGlobalVariables leaves the field poison in the init kernel:
store %struct.Agg { ptr addrspace(4) poison, [4 x i32] [...] }, ...
The fixture runs the real post-link pipeline with opt, so it needs no
device and runs on every lane and under both
CHIP_ENABLE_DEVICE_PROGRAM_SCOPE_GLOBALS settings. @s's initializer is
the one hipcc emits for
__device__ int Scalar = 42;
struct Agg { int *P; int A[4]; };
__device__ Agg S = {&Scalar, {1, 2, 3, 4}};
and the kernel adds the other way the same early return is reached: a
constant aggregate operand of an instruction, here a <2 x i64> of
ptrtoint addresses, which the unfixed pass stores as poison too. The
test fails when any store in the output has a poison operand; on the
unfixed tree both stores do.
A device-level test of the struct case was dropped in review in favour
of this one. On an Arc B570 it printed PtrOK=0 A3=4 before the fix and
PtrOK=1 A3=4 after it on Level Zero and OpenCL, but it could not run on
the OFF lane: there the fix leaves Scalar a program-scope global (its
address is loaded from S's init kernel, which isConvertibleGlobal
rejects), and rusticl rejects initialized program-scope globals (#1279),
so it needed a CMake gate and never exercised the vector path.
expandConstant rewrites a reference to a lowered device global into a
load of its __chip_var_<name> address holder, recursing through
constant expressions, but returned a ConstantAggregate (struct, array
or vector constant) unchanged. An initializer such as
__device__ int Scalar = 42;
__device__ Agg S = {&Scalar, {1, 2, 3, 4}};
takes path B of emitGlobalVarInitBody (hasNoRuntimeConstants already
recurses into aggregates and sees the reference), so the init kernel
stored the original constant, and when eraseMappedGlobalVariables
replaced @scalar with poison the field became poison:
store %struct.Agg { ptr addrspace(4) poison, [4 x i32] [...] }, ...
The device then read a wrong S.P. The same early return left an
aggregate operand of an ordinary kernel instruction referring to the
original global too; eraseMappedGlobalVariables marks that
llvm_unreachable, and a Release build poisons the operand instead (a
kernel storing <2 x i64> <ptrtoint @A, ptrtoint @b> stored poison).
Expand the operands like a constant expression's. If none changed the
aggregate is returned as before; otherwise it is rebuilt at the
insertion point with one insertvalue per operand, or insertelement for
a vector (a <2 x i64> of ptrtoint addresses reaches this path as an
instruction operand). IRBuilder folds the inserts of constant leading
operands, so a large aggregate costs instructions only from its first
runtime operand on. The init kernel for S now stores
%5 = insertvalue %struct.Agg poison, ptr addrspace(4) %4, 0
%6 = insertvalue %struct.Agg %5, [4 x i32] [i32 1, i32 2, i32 3, i32 4], 1
Rebuilding from poison was chosen over writing the non-runtime part with
path A's memcpy and patching the pointer fields through GEPs, which
would need a second copy of the initializer and per-field offsets.
Under CHIP_ENABLE_DEVICE_PROGRAM_SCOPE_GLOBALS=OFF the referenced
global is now loaded from another global's init kernel, so
lowerGlobalsToKernelArgs keeps it program-scope, exactly as it already
does for a plain `__device__ int *P = &Scalar;` initializer.
Fixes #1640
Collaborator
Author
|
/run-aurora-ci |
pvelesko
marked this pull request as ready for review
September 19, 2026 05:18
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.
expandConstant in HipGlobalVariables returned constant aggregates unchanged, so a lowered device global's address nested in a struct or array initializer (or in an aggregate operand of an instruction) became poison once the original global was erased. It now expands the aggregate's operands like a constant expression and, when any of them needs a runtime load, rebuilds the aggregate with insertvalue/insertelement. The reproducer runs the post-link pipeline with opt, so it covers both program-scope-globals settings without a device.
A non-lowered global initialized with a device global's address (a function-local static, or the constant behind a local aggregate) is a separate path, tracked in #1650.
Fixes #1640
This PR was generated using AI