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

[vyncint]

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

1. An arch-derived cfg for device codegen (cfg(cuda_oxide_sm_80) or similar), so warp_reduce can specialize where the target allows. This is the general fix and would serve every other arch-gated API in cuda-device, not just this one.
2. Recognise it in the compiler: match the butterfly in mir-lower and rewrite to redux.sync when the target is sm_80+. The backend does know the target. Pattern-matching a shuffle tree is fragile, though, and it would fight the convergent modelling.
3. Do nothing to the API, document the gap. Callers who know they are on Ampere+ and reducing integers can already write warp::redux_sync_add directly, 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.)

[vyncint]

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, and warp_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 below sm_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-device has the same problem. Happy to pick that up if you decide which shape you want.

[nihalpasham]

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-device API has the same hole.

Why not the others, plus one not listed:

• Matching the shuffle tree in mir-lower: fragile, and a walk-and-replace outside the conversion framework.
• A __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, sees redux.sync in 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 cfg is 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.

env (the backend reads these already)   floor   cfgs set (cumulative: cfg cannot compare numbers)
CUDA_OXIDE_TARGET=sm_86                 86      cuda_oxide_sm_at_least = "70" "75" "80" "86"
CUDA_OXIDE_DEVICE_ARCH=sm_120a (run)    120     ... "90" "100" "120"   (run, sanitize and debug set this hint)
neither                                 80      "70" "75" "80"   (today's cross-compile default)

Design I would recommend:

• Cfg: cuda_oxide_sm_at_least = "XY", one per floor, cumulative. No a suffixes for now.
• Set by a build.rs in cuda-device, from the env vars the backend already reads.
• The same floor goes into target selection as a minimum.
• warp_reduce: integer ops at N == 32 use redux.sync under the cfg. Everything else keeps the shuffle tree.
  • The N == 32 limit is scope, not an ISA rule: redux.sync reduces over whatever lanes are in membermask, 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 the 0xffffffff contract warp_reduce has; the sub-warp path has out-of-tile lane substitution that needs its own check before switching.
• Tests: PTX at sm_80 shows redux.sync, at sm_75 the tree. N = 16 and f32 as controls.
• One behavior change to state: a module that picks sm_75 today moves to sm_80 once it calls an integer warp_reduce.

[vyncint]

Taking this, building your option 1 as specced. Approach note first, since three things are judgment calls rather than transcription.

build.rs reuses cuda-target-spec rather than parsing arches again. That crate has no dependencies, so it is safe as a build-dependency of cuda-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. PipelineConfig documents device_arch_hint as "used only when no explicit target is provided", so build.rs resolves CUDA_OXIDE_TARGET > CUDA_OXIDE_DEVICE_ARCH > 80, which is your table exactly. It will also emit rerun-if-env-changed for both, and rustc-check-cfg so the new cfg does not trip unexpected_cfgs under -D warnings.

The fast path needs a trait hook, since Rust has no specialization. warp_reduce is generic over T and Op, so I plan a provided method on ReduceOp<T> -- returns None by default, overridden under cfg(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 exact redux_sync_* sibling; f32 keeps the butterfly as you said. warp_reduce consults it only when N == 32, so the sub-warp lane-substitution path is untouched.

On "the same floor goes into target selection as a minimum". Since build.rs reads 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 emits redux.sync gets 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 = 16 and f32 as controls, plus an A10G run for correctness and a re-measure of the 4.74x.

[vyncint]

Status: the cfg half is built and verified on hardware; the warp_reduce half 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:

env                                   cfgs set
(none)                                70 75 80
CUDA_OXIDE_TARGET=sm_75               70 75
CUDA_OXIDE_DEVICE_ARCH=sm_120a        70 75 80 86 90 100 120
TARGET=sm_86 + DEVICE_ARCH=sm_120     70 75 80 86    <- explicit target wins
CUDA_OXIDE_TARGET=""                  70 75 80       <- blank is unset, not a parse error

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_75 and run on the A10G, with all nine integer pairs, f32, and a WarpTile<16> in one kernel:

PTX @ sm_75: .target sm_75   redux.sync 0   shfl.sync 64
RUN @ sm_75: MATRIX PASS     BENCH us_per_launch=159.7

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:

Verification failed for lowered LLVM module:
  ["crates/cuda-device/src/cooperative_groups.rs": line 840] verification failed.
ReturnOp operand type does not match the function's result type
  Failed operation: llvm.return v2087

Four things it is not

Each probe below was run on the A10G at --arch sm_80 and each one passes, so the obvious explanations are eliminated:

probe	result
redux_sync_add called directly in a kernel (the redux_sum shape)	496
the same call through a plain non-generic #[inline(always)] helper	496
a bare trait method returning the result, dispatched through a generic fn dispatch<R: Reduce>	496
all three shapes again with shuffle_xor in place of redux_sync_add	pass

So 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 in return, and two whole #[cfg]-selected definitions each with a single-expression body matching WarpShuffle::shfl_xor_via exactly -- 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_FORM sitting in the same impl, and warp_reduce still 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, with enum payload storage type mismatch: si32 cannot be adapted to i32, for T = i32. An Option<i32> crossing a device function boundary is enough to trigger it. That is why the current shape uses a const plus 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]

@vyncint I dug into the lowering blocker only; your arch cfg and warp_reduce work stay with you.

It is narrower than the associated const or the dead butterfly branch. Any device function that returns a redux.sync result directly fails:

#[inline(never)]
fn reduce(v: u32) -> u32 { warp::redux_sync_add(FULL_MASK, v) }

convert_redux in mir-lower declared the intrinsic with the dialect op's unconverted ui32/si32 result while its operands were already signless (llvm_nvvm_redux_sync_add: ui32 (i32, i32)), so llvm.return fails 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-return hook shape builds at --arch sm_80 and emits redux.sync.add.s32 with no shfl. Your branch should build on top of it once it lands.

[nihalpasham]

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_reduce API specialization remain separate, so this issue stays partial.

[vyncint]

Unblocked by #1264 and up as #1305 — the target-derived cfg and the warp_reduce specialization, built as you specced in option 1.

At the Ampere floor the six u32 ops become six redux.sync with no shfl at all, against 30 shuffles at sm_75; the three signed ops go 3 against 15. The f32 and WarpTile<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_demo is 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_80 the module was rejected for ReduxF32 with no f32 redux in it, because the feature scan splits LLVM IR on ; and matched redux.sync and .f32 547 lines apart (#1303).

Leaving this issue open for you to close with the merge, since the 4.74x in the title is the primitive measurement from the original investigation rather than a claim about this PR.