Repository navigation
Conversation
NVIDIA#1305 selects the reduction form at compile time from the target floor, so the comparison has to be two builds of one source rather than two kernels in one binary. This adds that harness, plus the control the comparison needs. - `api_chain` / `api_independent` drive `warp_reduce::<u32, _>`: one is a dependency chain (reduction latency), the other four unrelated reductions per iteration (throughput). - `shuffle_chain` / `shuffle_independent` spell the butterfly out with `WarpCollective::shfl_xor`, so at `sm_80` the pair isolates the reduction from the rest of the build: same target, same compiler, same binary. - 128 threads, 40,000 iterations, 15 interleaved samples and 3 warm-up rounds per arm, medians reported per reduction, and the first iterations checked against a host simulation so a folded-away loop cannot pass as a timing. Measured on an A100-SXM4-40GB (cc 8.0, driver 595.91.07). At `sm_80`, where the API takes `redux.sync.add.s32` and the control keeps its five `shfl.sync.bfly.b32` + `add.s32`: | workload | redux ns/reduction | butterfly ns/reduction | ratio | |---|---|---|---| | dependent chain | 8.4 - 10.8 | 27.2 - 35.0 | 3.24x | | four independent reductions | 2.9 - 3.7 | 9.4 - 12.1 | 3.24 - 3.27x | Three batches agree on the ratio to the printed digit. Built at `sm_75` the two arms are the same code and time identically, which is the control for the harness itself. Needs NVIDIA#1305 for the `sm_80` arm to be a `redux.sync` build; without it both arms are the butterfly. Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>
Keeps the measurement next to the change it supports, so the numbers in the description can be checked against the transcript that produced them. Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>
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.
What
#1305 chooses the reduction form from the target floor at compile time, so the
two forms cannot be compared inside one binary — the comparison has to be two
builds of one source. That harness did not exist, and #1305 says so: "no new
speedup measurement was made."
Four arms, two per form:
api_chain/api_independentdrivewarp_reduce::<u32, _>. One is adependency chain, the other runs four unrelated reductions per iteration.
shuffle_chain/shuffle_independentspell the butterfly out withWarpCollective::shfl_xor, so on one target the pair isolates the reductionfrom the rest of the build: same compiler, same binary.
Correctness
The oracle is a per-warp host simulation of the real data flow, not a
comparison between the two arms — two arms can share a fault. Every reduction
reads the previous round's whole warp and only then broadcasts, so no lane
observes a value another lane already updated.
a, b, c, dinstead of folding them into one XOR, so a fault in one operation is locatableu32::MAX, the top bit,i32::MAX, one, fixed-seed values, cycled so different warps differMetric definitions
The number is an amortised collective cost, not instruction latency. It is
timed with a host wall clock around launch plus stream synchronisation, so
launch and host overhead are small per reduction but not absent, and one warp
reduction is one collective — not one instruction, and not 32 results.
Ris 1 for the chain arms and 4 for the independent ones and is printed perrow. Arms are sampled round-robin inside one pass, so a clock or power drift
lands on all of them. No threshold is asserted: the harness reports, it does not
gate.
A scaling check runs the chain arms at four iteration counts and warns when time
does not grow with the count — what a loop the compiler removed would look like.
Measured
A100-SXM4-40GB (cc 8.0, driver 595.91.07, no MIG), rustc
nightly-2026-08-28,llcLLVM23.1.0-rust-1.100.0-nightly, CUDA13.0.88.Test tree 1 —
mainalone (ec4aa479), no fast-path implementation:correctness passes, and
api_chainreports 27.19 ns, the same as thebutterfly. The harness measures the API as the fallback it actually compiled
to; it does not label the run
reduxbecause the target issm_80.Test tree 2 — #1305's dependency tree (
cf4c55c3) plus this harness,--arch sm_80:api_chainshuffle_chainapi_independentshuffle_independentPaired on one target and one binary: 3.23× for the dependency chain and
3.11× for the independent mix. Scaling,
api_chain/shuffle_chain:0.048/0.145, 0.221/0.704, 0.866/2.800, 1.727/5.595 ms at 1k, 5k, 20k, 40k
iterations.
Same dependency tree,
--arch sm_75: the control.api_chainandshuffle_chainboth report 27.19 ns, and the independent arms both 9.01 ns —the API compiles to the butterfly and the two arms are the same code.
predicate_kind=is not reported: the per-kernel PTX path proof from theprevious revision of this PR is kept outside the diff and is still available.
Reproducing
Needs #1305 for the
sm_80arm to be aredux.syncbuild; without it both armsare the butterfly and the harness still runs.
Limitations
timer_kindis a host wall clock, not CUDA events: this host API path wasverified against the pinned
cuda-coreand the event type is available, butthe harness keeps one timing path rather than two. The figure is bounded by
that, which is why it is named an amortised cost.
cargo oxide inspectstops at PTX and the driver JITs it atload, so path evidence is PTX.
sm_75is compile and runtime coverage on an A100 host, not a Turing GPU run.Scope
Measurement harness and documentation. It does not touch
crates/cuda-device,does not change the floor rule, and does not restate #1305's implementation,
which stays with its author and its #1304 dependency. Raw experiment artifacts
are kept outside this diff.