Repository navigation
warp_reduce runs a 10-instruction shuffle tree where Ampere's redux.sync does it in one: 4.74x measured, but no way to gate it on the target #811
Description
cooperative_groups::warp_reduce<T, Op, N> is a generic butterfly: log2(N) rounds of shfl_xor_via + Op::combine. For a full warp that is 5 shuffles + 5 combines = 10 instructions. redux.sync.add.u32 does the same job in one, cuda-oxide already exposes it as warp::redux_sync_add (#223, #258) — the generic API just never routes integers to it.
Measured
A10G (sm_86, driver 595.71.05), 1M threads x 64 full-warp u32 sum reductions each, 30 timed launches after 5 warmup. Same kernel either way; only the reduction call differs:
| form | us/launch | correct |
|---|---|---|
warp_reduce::<u32, Sum, _>(&tile, v) |
159.4 | yes |
warp::redux_sync_add(FULL_MASK, v) |
33.6 | yes |
| 4.74x |
Both bit-identical against a host replay, which is expected: wrapping u32 addition is associative mod 2^32, and min/max/and/or/xor are associative and exact, so the reduction order cannot change the result.
PTX confirms the mechanism — 10 shfl.sync versus 8 redux.sync across the respective unroll factors, i.e. 10 instructions per reduction against 1.
Scope of that number: the probe kernel does 64 reductions and almost nothing else, so it is reduction-bound and this is the primitive's cost in isolation, not an application speedup. It sizes the prize; a memory-bound kernel with one reduction at the end would see little.
The participation contract already matches
This looked like it might be unsound, and it is not. #672 states the existing behaviour plainly: "The warp reductions shuffle with the full 32-lane member mask, so every lane must be launched and converged." #672 then added separate _partial entry points for short warps rather than relaxing warp_reduce.
So warp_reduce on a WarpTile<32> already requires a fully converged warp and the full mask — exactly redux.sync's requirement. Substituting it does not weaken anything. Sub-warp tiles (N < 32) would keep the shuffle path, since each tile carries its own mask and the wrapper's out-of-tile lane substitution is shuffle-specific.
What actually blocks it
redux.sync is sm_80+ and the PTX will not assemble below that. redux_sum's own comment says so, and that example handles it by checking ctx.compute_capability() at runtime and printing skipping: redux.sync.add requires sm_80+ (Ampere).
Device code has no way to ask what it is being compiled for. There is no arch-derived cfg supplied to device codegen — I looked; the sm_80+ requirement lives in doc comments and is enforced by the caller. --device-cfg exists but is user-supplied, not automatic.
So warp_reduce cannot simply emit redux.sync: every sm_75 build that uses a warp reduction would stop assembling. That is a hard regression, not a tuning question, which is why this is an issue and not a PR.
Options, all yours to pick
- An arch-derived cfg for device codegen (
cfg(cuda_oxide_sm_80)or similar), sowarp_reducecan specialize where the target allows. This is the general fix and would serve every other arch-gated API incuda-device, not just this one. - Recognise it in the compiler: match the butterfly in mir-lower and rewrite to
redux.syncwhen the target is sm_80+. The backend does know the target. Pattern-matching a shuffle tree is fragile, though, and it would fight theconvergentmodelling. - Do nothing to the API, document the gap. Callers who know they are on Ampere+ and reducing integers can already write
warp::redux_sync_adddirectly, and a measured note in the book would tell them it is worth 4.74x on the primitive. Cheapest, and no risk.
Happy to implement whichever you prefer, including (3) on its own. I did not want to pick the mechanism for you when the enabling piece — target-aware device code — is a broader design decision than this one reduction.
(Related but distinct: #274 tracks adding SM80+ intrinsics. redux.sync is already added; this is about the high-level API not using it.)
Activity
vyncint commentedon Aug 12, 2026
Took option 3 as a first step: #813 documents the gap without touching the API.
The subsection under the book's "Warp reduction" names all eight
redux_sync_*functions with the measurement, andwarp_reduce's own rustdoc carries a short note, since that is what someone about to call it reads. Both state the scope of the 4.74x plainly — it sizes the primitive, not an application — and both spell out the three constraints (integers only, full converged warp, will not assemble belowsm_80).Leaving this issue open. #813 deliberately does not close it: options 1 and 2 are still the real fix, and the enabling piece — device code knowing its target architecture — is worth more than this one reduction, since every other arch-gated API in
cuda-devicehas the same problem. Happy to pick that up if you decide which shape you want.docs(warp): point integer reductions at redux.sync, with the measurem…
nihalpasham commentedon Sep 3, 2026
We want this, and option 1 (an arch-derived cfg) is the one I would take. Device code that knows its floor (the lowest target the PTX will be built for) is better; every arch-gated
cuda-deviceAPI has the same hole.Why not the others, plus one not listed:
__CUDA_ARCH__-style intrinsic folded at lowering: the requirement scan (the pass that reads which instructions a module uses to pick its target) runs on typed IR before lowering, seesredux.syncin the dead branch, and raises the floor to sm_80 before anything folds. A cfg avoids this because rustc drops the dead branch before MIR exists.A
cfgis decided before rustc runs, so it can only carry a lower bound. That matches how targets are picked: the backend honors the floor whenever the module can run on it, and only moves for modules that could not run there anyway.Design I would recommend:
cuda_oxide_sm_at_least = "XY", one per floor, cumulative. Noasuffixes for now.build.rsincuda-device, from the env vars the backend already reads.warp_reduce: integer ops at N == 32 useredux.syncunder the cfg. Everything else keeps the shuffle tree.redux.syncreduces over whatever lanes are inmembermask, so a sub-warp tile could use it with its own mask. Full warp first because it is the measured case and its mask is already the0xffffffffcontractwarp_reducehas; the sub-warp path has out-of-tile lane substitution that needs its own check before switching.redux.sync, at sm_75 the tree. N = 16 and f32 as controls.warp_reduce.vyncint commentedon Sep 5, 2026
Taking this, building your option 1 as specced. Approach note first, since three things are judgment calls rather than transcription.
build.rsreusescuda-target-specrather than parsing arches again. That crate has no dependencies, so it is safe as a build-dependency ofcuda-device, and it keeps one arch parser in the tree instead of a second one that can disagree with the backend.Precedence is copied from the backend, not invented.
PipelineConfigdocumentsdevice_arch_hintas "used only when no explicit target is provided", sobuild.rsresolvesCUDA_OXIDE_TARGET>CUDA_OXIDE_DEVICE_ARCH> 80, which is your table exactly. It will also emitrerun-if-env-changedfor both, andrustc-check-cfgso the new cfg does not tripunexpected_cfgsunder-D warnings.The fast path needs a trait hook, since Rust has no specialization.
warp_reduceis generic overTandOp, so I plan a provided method onReduceOp<T>-- returnsNoneby default, overridden undercfg(cuda_oxide_sm_at_least = "80")for the nine integer impls (u32: Sum/Min/Max/BitAnd/BitOr/BitXor,i32: Sum/Min/Max). All nine have an exactredux_sync_*sibling;f32keeps the butterfly as you said.warp_reduceconsults it only whenN == 32, so the sub-warp lane-substitution path is untouched.On "the same floor goes into target selection as a minimum". Since
build.rsreads the same two env vars with the same precedence, the cfg floor and the selected target already agree in all three rows of your table, and a module that emitsredux.syncgets raised to sm_80 by the requirement scan regardless. So rather than add clamping logic in mir-importer, I plan to pin the agreement with a test over the three cases and say so in the PR. If you would rather have an explicit clamp in target selection as belt-and-braces, say so and I will add it -- I did not want to widen the blast radius into the backend on my own reading.One consequence worth stating up front, since it is the behaviour change you flagged: the no-env floor is 80, so this turns on by default for cross-compiles, and a module that would have picked sm_75 moves to sm_80 the moment it calls an integer
warp_reduce.Evidence plan: PTX at sm_80 vs sm_75 from
coop_groups_demo(the only example on this path),N = 16andf32as controls, plus an A10G run for correctness and a re-measure of the 4.74x.vyncint commentedon Sep 5, 2026
Status: the cfg half is built and verified on hardware; the
warp_reducehalf is blocked on a lowering failure I have narrowed but not explained. Reporting now rather than sitting on it, since the remaining question looks like it is in your half of the tree.What works, verified on an A10G (sm_86, driver 595.71.05)
Floor resolution and the cfg ladder behave exactly as your table specifies, measured by reading the cfgs back from inside the crate:
The rule lives in
cuda-target-spec(resolve_sm_floor,SM_FLOOR_LADDER,sm_floors_at_most) so the backend and the build script share one implementation, and a named target that does not parse fails the build rather than falling back.Below the floor it is correct end to end. Built at
--arch sm_75and run on the A10G, with all nine integer pairs,f32, and aWarpTile<16>in one kernel:159.7 us matches the 159.4 us butterfly figure from the original measurement, so the baseline reproduces.
What is blocked
At
--arch sm_80, where the hook turns on, device codegen rejects the module:Four things it is not
Each probe below was run on the A10G at
--arch sm_80and each one passes, so the obvious explanations are eliminated:redux_sync_addcalled directly in a kernel (theredux_sumshape)#[inline(always)]helperfn dispatch<R: Reduce>shuffle_xorin place ofredux_sync_addSo it is not indirection, not a trait, not generic dispatch, and not redux-through-a-wrapper as such. I also tried two body shapes for the hook -- a
#[cfg]block ending inreturn, and two whole#[cfg]-selected definitions each with a single-expression body matchingWarpShuffle::shfl_xor_viaexactly -- and both fail identically, so it is not the diverging-block shape either.What remains different from the passing repro: the associated
const HAS_FULL_WARP_FORMsitting in the same impl, andwarp_reducestill carrying the dead butterfly branch after the early return. I have not isolated which.Worth recording separately, found on the way: returning
Option<T>from this hook fails earlier and differently, withenum payload storage type mismatch: si32 cannot be adapted to i32, forT = i32. AnOption<i32>crossing a device function boundary is enough to trigger it. That is why the current shape uses aconstplus a plain return.Where I would like your steer
The branch is complete and green locally at both floors (106 tests each, clippy
-D warnings, fmt, rustdoc, and the book), and I am happy to keep bisecting. But the failure is in lowering rather than in the device API, so you may recognise it at a glance or prefer to take it. Say which and I will either keep going or hand over the reproducers.No PR yet: the fast path is the point of the change and it does not compile at the floor it targets, so there is nothing worth reviewing.
sachinsharma3191 commentedon Sep 13, 2026
@vyncint I dug into the lowering blocker only; your arch cfg and
warp_reducework stay with you.It is narrower than the associated const or the dead butterfly branch. Any device function that returns a
redux.syncresult directly fails:convert_reduxin mir-lower declared the intrinsic with the dialect op's unconvertedui32/si32result while its operands were already signless (llvm_nvvm_redux_sync_add: ui32 (i32, i32)), sollvm.returnfails verification. A result that only feeds a store never hits this, which is why the in-kernel probes passed.The fix and regression tests are in #1264. With it, the
const HAS_FULL-plus-returnhook shape builds at--arch sm_80and emitsredux.sync.add.s32with noshfl. Your branch should build on top of it once it lands.fix(mir-lower): lower redux.sync results to signless integers
nihalpasham commentedon Sep 18, 2026
I reviewed #1264 and added controls for all reduction-lowering routes. Out-of-line redux helpers now return the translated LLVM result type, fixing the return-verification failure.
Native and modern libNVVM validation pass locally. The legacy unresolved-redux limitation also reproduces without this change. Hosted CI also passes on this head.
This addresses one compiler-side prerequisite. The target-derived cfg and
warp_reduceAPI specialization remain separate, so this issue stays partial.vyncint commentedon Sep 19, 2026
Unblocked by #1264 and up as #1305 — the target-derived cfg and the
warp_reducespecialization, built as you specced in option 1.At the Ampere floor the six
u32ops become sixredux.syncwith noshflat all, against 30 shuffles at sm_75; the three signed ops go 3 against 15. Thef32andWarpTile<16>controls are byte-identical at both floors. Correct on an A10G at the auto-detected arch and forced to sm_75, smoketest 3/3.Two notes. It needed a new example:
coop_groups_demois the only caller of the generic API and it also uses clusters, so it cannot build below sm_90 — there was nowhere to read the difference. And it is stacked on #1304: at--arch sm_80the module was rejected forReduxF32with no f32 redux in it, because the feature scan splits LLVM IR on;and matchedredux.syncand.f32547 lines apart (#1303).Leaving this issue open for you to close with the merge, since the
4.74xin the title is the primitive measurement from the original investigation rather than a claim about this PR.fix(codegen): match redux.sync's f32 form on the instruction, not on …
docs(evidence): A100/SM80 execution record for the 2026-09-26 pack
docs(evidence): A100/SM80 execution record for the 2026-09-26 pack
fix(codegen): match redux.sync's f32 form on the instruction, not on …