Parallelize program-scope variable init, fill large zero ones from host (#582) - #1379
Open
pvelesko wants to merge 1 commit into
Open
Parallelize program-scope variable init, fill large zero ones from host (#582)#1379pvelesko wants to merge 1 commit into
pvelesko wants to merge 1 commit into
Conversation
…st (#582) PR #1355 merged the per-variable init shadow kernels into a single __chip_var_init_all, removing the O(N) launch overhead, but left the copy itself running on one work item. That single work item is what issue #582 actually complains about: a 1 MiB variable took 32.7 ms to initialize on an Arc B570 regardless of how few launches it took to get there. Three changes, all in the initialization path: Non-zero initializers are now copied by a grid-stride loop driven by get_global_id(0)/get_global_size(0) instead of a single-work-item llvm.memcpy, and the runtime launches the init kernel with a real geometry sized from the bytes it already knows about. The element type widens to i32/i64 when size and alignment allow, since byte granularity wastes most of the available bandwidth. Zero initializers below ChipVarFillThreshold become a grid-stride store loop that reads no initializer blob at all. Above the threshold they leave the kernel entirely and the runtime zeroes the storage with memFillAsync, which reaches a DMA engine. The threshold exists because host fills do not amortize: 256 eight-byte variables cost ~1782 us as individual fills versus ~99 us as one combined kernel, so an unconditional switch to fills would be a large regression for the many-small-variables case. CHIPVarInfo[2] carries which of the three cases applies. It stays a tri-state in the existing slot rather than growing the array, because the HIPRTC on-disk cache keys SPIR-V on nothing derived from the pass plugin: a module built by an older pass can be handed to a newer runtime, and a tri-state degrades correctly to "initialized by the kernel" where a fourth slot would be read out of uninitialized device memory. Two latent problems the new geometry would otherwise have exposed: - Initializers that reference other lowered variables are a single scalar store, not a partitioned loop, so with a real grid every work item executed the same store to the same address. Now guarded to one work item. - QueuedKernels gated the only Queue->finish(), which a module initialized entirely by host fills would have skipped. Initialization runs on the default queue while the consuming kernel may run on another stream, so that barrier is load-bearing. Measured on Intel Arc B570, first-use init, median of 9, level0 / opencl: 4 MiB non-zero 72.1 -> 43.2 ms / 83.1 -> 50.1 ms 4 MiB zero 37.9 -> 20.3 ms / 38.1 -> 17.4 ms 1 MiB zero 24.2 -> 20.6 ms / 23.6 -> 17.8 ms 512 KiB non-zero 23.5 -> 20.6 ms / 24.3 -> 20.6 ms 256 scalars unchanged (dominated by info/bind kernel launches) check.py passes 100% on both backends (1045 level0, 1013 opencl). Refs #582. Initializing non-zero variables by host-side copy is still not possible without a new channel for the initializer bytes: hip-clang emits the host shadow variables as undef, so the data exists only inside the device module.
pvelesko
force-pushed
the
fix-582-dma-init
branch
from
July 28, 2026 13:01
22356d6 to
fb98821
Compare
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.
Follow-up to #1355. That PR merged the per-variable init shadow kernels into one
__chip_var_init_all, removing the O(N) launch overhead, but left the copy itself on a single work item, which is what #582 actually asks about.Non-zero initializers are now copied by a grid-stride loop (get_global_id/get_global_size) with the element type widened to i32/i64 by size and alignment, and the runtime launches the init kernel with a real geometry. Zero initializers below a threshold become a store loop that reads no initializer blob; above it they leave the kernel and the runtime zeroes the storage with memFillAsync (DMA). The threshold exists because host fills do not amortize: 256 eight-byte variables cost ~1782us as individual fills versus ~99us in one kernel.
CHIPVarInfo[2] carries which case applies as a tri-state in the existing slot rather than a fourth slot, because the HIPRTC on-disk cache keys SPIR-V on nothing derived from the pass plugin, so an older module can reach a newer runtime; a tri-state degrades to "initialized by the kernel" where a fourth slot would read uninitialized device memory.
Also fixes two problems the new geometry would have exposed: initializers referencing other lowered variables are a single scalar store and are now guarded to one work item, and the only
Queue->finish()was gated on a kernel having been queued, which a fill-only module would have skipped.Measured on Intel Arc B570, first-use init, median of 9, level0 / opencl:
256 scalars is unchanged because that case is dominated by the 256 info and bind kernel launches, which this does not touch.
check.py: 100% pass on both backends (1045 level0, 1013 opencl). New test TestGlobalVarInitClasses covers all six initializer classes in one module.
Refs #582. Not closing it: initializing non-zero variables by host copy still needs a channel for the initializer bytes, since hip-clang emits the host shadow variables as undef and the data exists only inside the device module.