Repository navigation
Conversation
A partial result of a reduction or scan whose operator has no neutral
element is an `Optional{T}`, the value and a `Bool` flag. For a bits
type the flag is padded to `T`'s alignment, and storing an `Optional`
writes the fields only, so the padding of partials in device memory was
never written. The second pass loads partials whole and the host copies
the last one, so Compute Sanitizer's initcheck reported every such
reduction, including Base reductions with a custom operator that
GPUArrays sends to AK.
The flag is now the unsigned integer that fills those bytes
(`Optional{T, F}` with `F = flag_type(T)`), up to 64 bits, which is as
wide as integers go in OpenCL's SPIR-V. Sizes are unchanged, and storing
an `Optional` writes all of it, unless `T` has padding of its own or is
aligned to more than 8 bytes. Other types keep a `Bool`.
maleadt
force-pushed
the
ab/optional-padding
branch
from
October 28, 2026 09:55
d00506f to
fb353af
Compare
This was referenced Oct 28, 2026
This branch has not been deployed
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
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.
A reduction or scan whose operator has no known neutral element (no
neutralkeyword and noGPUArraysCore.neutral_elementmethod) keeps its partial results asOptional{T}: the value and aBoolflag, padded toT's alignment. LLVM stores the fields separately (st.global.b64+st.global.b8), but the second pass loads partials whole (ld.global.nc.v2.b64) and the host copies the last one whole, reading bytes nothing wrote. Results are correct, but Compute Sanitizer's initcheck reports every such reduction, including Basereducewith a custom operator, which GPUArrays 12 sends to AK:This PR makes the flag the unsigned integer that fills those bytes:
Optional{T, F}withF = flag_type(T), e.g.UInt64forFloat64. Sizes don't change, and storing anOptionalwrites all of its bytes. Bits unions and non-bits types keep aBool. CUB's scan tile descriptors do the same. The flag is at most 64 bits, as OpenCL's SPIR-V has no wider integers, so 16-byte aligned types keep 8 bytes of padding; padding insideTisn't covered either. Storing the partial whole doesn't help, since LLVM treats padding as undefined.A script of reductions and scans without a neutral element gives 309 initcheck reports on main and 2 with this PR (from
Int128andNTuple{4, VecElement{Float32}}). The test checks that partials have no padding for common element types (failing on main) and that the flag stays within 64 bits. reduce/accumulate tests pass on CUDA (RTX 5080, Orin, Xavier, Jetson Nano), AMDGPU, PoCL, oneAPI and Metal.The layout changes codegen. Measured on top of #194, which keeps
Optionalpartials out of stack slots (without it, the comparison is dominated by ptxas promoting those slots only for the wide flag), the time ratios of this PR forFloat32max, 10^4 to 10^8 elements, median of 3 runs:accumulate!accumulate,dims=1On PoCL, the work-group vectorizer accesses the wider flag in local-memory scan tiles with AVX-512 gathers and scatters, where it leaves the
i8flag scalar. The Metal scans alongdimsalso keep the flag in threadgroup memory, now asi32accesses instead ofi8; I couldn't see why that's slower without the AGX assembly. Choosing the flag type per backend through the tunings was considered: it takes 15–20 lines throughreduce_seedand the kernel-side constructors, and loses initcheck coverage on those backends, so I left it out.