feat(cuda-device): derive an SM floor cfg and use redux.sync in warp_reduce by vyncint · Pull Request #1305 · NVIDIA/cuda-rust

vyncint · 2026-09-19T01:51:10Z

What this adds

Full-warp integer warp_reduce calls use redux.sync when the target floor is sm_80 or newer. Floats, smaller tiles and lower floors retain the shuffle implementation.

full-warp integer reduction at sm_80 -> one redux instruction
same reduction at sm_75              -> shuffle tree

The device build script and both backend selectors share the floor rule: an explicit target sets it; an advisory device hint can only lower the default sm_80 floor. Automatic fallback targets are raised to that floor or rejected if their feature requirements conflict. A numeric floor does not promise architecture-family-specific instructions.

Verification

At 96ea0b56, 1,395 tests/doctests across cuda-device, cuda-target-spec and cuda-oxide-codegen pass from an isolated build directory, along with strict Clippy and formatting. The new malformed-hint regression checks all three stages: the floor resolver and PTX selector ignore the hint; NVVM rejects it.

Earlier B200 runs at db5c23f8 passed for sm_80 and sm_75 through LLVM NVPTX, covering nine integer operation/type pairs plus float and sub-warp controls. The sm_80 run also passed memcheck and synccheck with 0 errors. The floor/reduction implementation and maintainer documentation fixes are unchanged in the new stack. GPU tests were not repeated after the scanner rebase; no new speedup measurement was made.

Remaining prerequisite

The original four #1304 lexical failures are fixed. Its updated PTX scan still misses valid whitespace/comment-separated opcode modifiers and treats a quoted .file path as an instruction. Those cases must be fixed before this PR lands. The earlier floor-selection counterexamples remain resolved.

Partially addresses #811. Depends on #1304.

nihalpasham

The integer reduction fast path is useful, but the cfg floor still needs to be part of target selection before this lands.

resolve_sm_floor is only called by cuda-device/build.rs. With no environment target, it advertises 80 while an sm_75-feature module selects sm_75. With a detected sm_120 hint and WGMMA, it advertises 120 while the backend selects sm_90a. I reproduced both through the actual resolver; the lower-bound assertion fails.

Please carry one resolved floor into native PTX and NVVM selection, handle incompatible architecture families explicitly, and add cross-stage regressions for these cases and invalid/blank inputs. The current tests only verify the new helper's own table.

This partially addresses #811; #1304 remains a prerequisite.

vyncint · 2026-09-23T07:33:15Z

You were right that the lower bound did not hold, and chasing it turned up that the rule itself — the middle row of the table on #811 — is the unsound part, not just my wiring of it.

Why a detected device cannot raise the floor

CUDA_OXIDE_DEVICE_ARCH is advisory by contract, said three times in the tree: "the backend builds for the detected GPU only when that GPU can actually run the kernel", and otherwise for the arch the kernel requires (apply_device_arch_hint, detect_run_target_arch, and the fallback in resolve_ptx_target_with_generated). A lower bound taken from a value the backend is documented to override is not a lower bound, which is exactly your second case: detected sm_120 with WGMMA advertises 120 and lands on sm_90a.

So the rule is now asymmetric, and that is what makes the assertion hold:

	floor
explicit target	that target, exactly — it is what gets built
detected device	lowers the floor, never raises it
neither	`DEFAULT_SM_FLOOR` (80)

The hint says nothing about how high selection goes. It does bound how low, because a device that can run the kernel is selected as-is. That direction matters: cargo oxide run on a Turing GPU must keep producing sm_75, and a blanket clamp to 80 would have built PTX that GPU cannot load. Pulling the floor down instead keeps device code honest about what it will get.

Carried into both selectors

Both resolve_ptx_target_with_generated and resolve_nvvm_target_with_generated now resolve the floor through the same resolve_sm_floor the build script uses — not a second copy of the rule — and raise a feature-selected target that lands below it. A floor that cannot host the module is reported rather than resolved to something else; nothing reaches that arm today, and it is there so a future floor cannot pick an architecture family silently.

Your two cases, now:

no env, ldmatrix-only module     PTX sm_80   NVVM sm_80    (was sm_75)
hint sm_120a, WGMMA module       PTX sm_90a  NVVM sm_90a   floor 80, so no contradiction
hint sm_75, basic module         PTX sm_75   (from "detected GPU"), floor 75

The regression you asked for is the property itself, not those three rows: a matrix over explicit × hint × features asserting selected.capability() >= resolve_sm_floor(explicit, hint) through both resolvers. Plus the blank and malformed inputs, where the two stages deliberately differ — a blank CUDA_OXIDE_TARGET is unset to the floor rule and a parse error to the backend, and the strict side is the one that decides whether anything is built; a malformed hint is ignored by both, matching the .parse().ok() filter in the PTX path.

One behaviour change, stated because it reaches builds that never use this cfg

With nothing pinned, a module whose features ask only for sm_75 is now built for sm_80. That is the documented Ampere+ minimum and the same default a featureless module already got, and --arch sm_75 still selects sm_75. It changed one existing test, which pinned the old semantics; I split it so it still asserts the feature's own floor is sm_75 and separately records the selected target.

Hardware

A10G, compute capability 8.6. Per-kernel PTX at both floors, unchanged from the last revision:

kernel	sm_80	sm_75
`reduce_u32_all_ops`	redux.sync 6, shfl 0	redux 0, shfl 30
`reduce_i32_signed_ops`	redux.sync 3, shfl 0	redux 0, shfl 15
`reduce_f32_control`	redux 0, shfl 15	redux 0, shfl 15
`reduce_subwarp_control`	redux 0, shfl 4	redux 0, shfl 4

Run on the A10G with nothing pinned (hint sm_86, floor 80, fast path live): all four checks correct. generated_ldmatrix unpinned now reports .target sm_80, the behaviour change above. Smoketest 4/4 over warp_reduce_redux, redux_minmax, redux_f32, generated_ldmatrix.

1218 codegen tests, plus cuda-device and cuda-target-spec; clippy -D warnings clean on all three; fmt; all 17 guards. Still stacked on #1304 — review the second and third commits.

nihalpasham

The earlier floor-contract blocker is fixed. Both selectors now use the shared floor: no-env ldmatrix stays at least sm_80, and an sm_120 hint with WGMMA advertises floor 80 while selecting sm_90a. I shortened the comments and corrected the outdated cfg and reduction-hook descriptions.

At db5c23f8, all 1,394 tests/doctests across cuda-device, cuda-target-spec and cuda-oxide-codegen passed, along with strict Clippy and formatting. B200 runs of warp_reduce_redux passed for sm_80 and sm_75 through LLVM NVPTX. The sm_80 fast path also passed memcheck and synccheck with 0 errors; float and sub-warp controls passed in both builds. No new performance measurement was made.

This still depends on #1304, whose current detector misses valid float-redux asm and misreads comments/ordinary symbols. Please fix those lexical cases there and update this stack before landing. The remaining request is the scanner prerequisite, not the repaired floor invariant.

One minor distinction: malformed hints are ignored by the shared floor resolver and PTX selector, but the existing NVVM selector rejects them. The floor guarantee is still fail-closed.

vyncint · 2026-09-24T13:48:41Z

Thanks for the B200 runs and the memcheck/synccheck pass — that is evidence I could not have produced on an A10G.

Rebased onto the updated #1304 (279ffe0f). Your docs: clarify the shared SM floor contract is carried through unchanged: all five files it touches are the same blobs as in db5c23f8, checked by object hash rather than by reading the diff.

On the minor distinction you noted: you are right, and I had it wrong in two places. My last comment here said a malformed hint "is ignored by both", and the test comment said "both stages ignore it" — true of the floor resolver and the PTX selector, not of the NVVM selector, which rejects it. One small commit on top, 96ea0b56, makes the comment name all three readings and asserts the NVVM rejection beside the PTX fall-through the test already checked, so the difference is a stated contract rather than something a later change could erase unnoticed. No behaviour changes; as you say, both outcomes are fail-closed.

Across the stack: cuda-target-spec 15, cuda-device 114, cuda-oxide-codegen 1266 — 1395, one more than your 1394 at db5c23f8, which is the one test function #1304 gained. Clippy -D warnings clean on all three, fmt clean, all 17 guards pass.

Stack is now 279ffe0f (#1304) ← 6dcbf00d ← df4ac829 ← c87f5f7e (yours) ← 96ea0b56.

nihalpasham

Checked 96ea0b56: the floor logic, reduction hooks and all five documentation-cleanup files are unchanged from the tested db5c23f8. The new test correctly pins NVVM's rejection of malformed hints beside PTX's fall-through.

All 1,395 tests/doctests across cuda-device, cuda-target-spec and cuda-oxide-codegen pass from an isolated build directory, along with strict Clippy and formatting. The earlier B200 sm_80/sm_75 runs and fast-path memcheck/synccheck covered the unchanged reduction implementation; they were not repeated for this scanner rebase.

The original four #1304 cases are fixed, but its PTX scan still misses legal whitespace/comment-separated modifiers and mistakes a quoted .file path for an instruction. Those cases are accepted by LLVM and ptxas, and the actual detector chooses the wrong targets. Please address that remaining prerequisite and update this stack. The floor invariant itself remains clear.

vyncint · 2026-09-24T14:38:59Z

Rebased onto #1304's 6b4ac7d0. The rebase changed nothing in this stack except #1304's two files — git diff 96ea0b56 cf4c55c3 touches only detect.rs and its test — and your five cleanup files are still the blobs you tested at db5c23f8.

cuda-target-spec 15, cuda-device 114, cuda-oxide-codegen 1267: 1396, your 1395 plus the one test #1304 gained. Clippy -D warnings clean on all three, fmt clean.

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>

Results for the four items in the cuda-oxide A100 execution pack (NVIDIA#399, NVIDIA#811 / NVIDIA#1305, NVIDIA#1235, NVIDIA#1328), with the raw logs behind every number. - A1: NVIDIA#399 closure audit (`CLOSURE_COMMIT: UNKNOWN`) and the test-only regression that covers the iterator form and its negative control. - A2: native A100 validation of NVIDIA#1305 at cf4c55c - correctness, sanitizers, sm_75 fallback, and the paired measurement (3.24x for a dependent chain, 3.24-3.27x for four independent reductions). - A3: NVIDIA#1235 merged-state and coverage audit. Existing oracle covers 2 of 12 combinations; G1/G2/G3 unmet, so no transform is proposed. - A4: NVIDIA#1328 - the reported reproducer is already fixed by NVIDIA#1264, the equal-width arm was still wrong, falsification and device evidence. Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

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>

Results for the four items in the cuda-oxide A100 execution pack (NVIDIA#399, NVIDIA#811 / NVIDIA#1305, NVIDIA#1235, NVIDIA#1328), with the raw logs behind every number. - A1: NVIDIA#399 closure audit (`CLOSURE_COMMIT: UNKNOWN`) and the test-only regression that covers the iterator form and its negative control. - A2: native A100 validation of NVIDIA#1305 at cf4c55c - correctness, sanitizers, sm_75 fallback, and the paired measurement (3.24x for a dependent chain, 3.24-3.27x for four independent reductions). - A3: NVIDIA#1235 merged-state and coverage audit. Existing oracle covers 2 of 12 combinations; G1/G2/G3 unmet, so no transform is proposed. - A4: NVIDIA#1328 - the reported reproducer is already fixed by NVIDIA#1264, the equal-width arm was still wrong, falsification and device evidence. Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

NVIDIA#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. This adds that harness and the control it needs. - `api_chain` / `api_independent` drive `warp_reduce::<u32, _>`. One is a dependency chain, the other four unrelated reductions per iteration. - `shuffle_chain` / `shuffle_independent` spell the butterfly out with `WarpCollective::shfl_xor`, so on one target the pair isolates the reduction from the rest of the build: same compiler, same binary. - Correctness is a per-warp host simulation of the real data flow, not a comparison between the two arms. Every reduction reads the previous round's whole warp and only then broadcasts, so no lane observes a value another lane already updated. Both arms are checked against that oracle at iterations 0, 1, 2, 3, 8 and 17, on every lane and every field. - The independent arms publish four values per lane instead of folding them into one, so a fault in a single operation is locatable. - Every arm starts from the bitwise inverse of its oracle, so a value the kernel never writes cannot read back as a correct answer. On a tree without NVIDIA#1305 both arms are the butterfly and the harness still runs: the API is then measured as the fallback it compiles to, not as a fast path. Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de> The figure the harness reports is a warp reduction's amortised cost over a long kernel, timed with a host wall clock around launch plus stream synchronisation. That keeps launch and host overhead small per reduction without removing them, and one warp reduction is one collective - not one instruction, and not 32 independent results. The output now says so: ``` timer_kind=host_wall_clock_kernel_batch iterations=40000 grid_blocks=1 warps_per_block=4 amortized_ns_per_warp_reduction = kernel_ms * 1e6 / (iterations * blocks * warps * R) ``` `R` is 1 for the chain arms and 4 for the independent ones, and it is printed per row. The arms are sampled round-robin inside one pass so a clock or power drift lands on all of them. A scaling check runs the chain arms at four iteration counts and warns when the time does not grow with the iteration count, which is what a loop the compiler removed would look like. On the dependency tree the four points are 0.048/0.145, 0.221/0.704, 0.866/2.800 and 1.727/5.595 ms. No performance threshold is asserted: the harness reports, it does not gate. Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

A reusable harness for the pair NVIDIA#1305 selects at compile time. Four arms: API and hand-written butterfly, each as a dependency chain and as four unrelated reductions per iteration. Correctness is a per-warp host simulation of the real data flow, at iterations 0, 1, 2, 3, 8 and 17, on every lane and field; the independent arms publish four values per lane instead of one merge, and every arm starts from the bitwise inverse of its oracle. Timing is reported as `amortized_ns_per_warp_reduction` with `timer_kind` and `R` stated, and a scaling check warns when the time does not grow with the iteration count. On NVIDIA#1305's dependency tree at sm_80: 8.41 ns against 27.19 ns for the chain, and 2.90 against 9.01 ns for the independent mix. On a tree without NVIDIA#1305 both arms are the butterfly and the harness reports the fallback rather than a fast path. Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>

…nearby text `contains_redux_f32_features` asked whether `redux.sync` and `.f32` both appeared in one `;`-delimited piece of the LLVM IR. `;` starts a comment in IR, so that piece is usually the whole module: a module reducing integers with any unrelated `.f32` was pushed from sm_80 to sm_100a. The integer `warp_reduce` fast path in NVIDIA#811 is such a module. The predicate now decides on tokens. The IR is lexed with quoted text consumed whole and comments dropped, and two things count: * a global symbol, quoting decoded, named `llvm.nvvm.redux.sync.` followed by a float operation -- the ones beginning with `f`, since the float callees carry no type suffix; * a `redux` opcode whose modifiers, in an inline-asm template, start with `.sync` and include `.f32`. The template is the string after `LLParser`'s `asm [sideeffect] [alignstack] [inteldialect] [unwind]`, decoded as `llvm::UnEscapeLexed` does, then lexed as PTX: whitespace and `//` and `/* */` comments separate tokens without ending a modifier run, a number consumes its whole run, and a string literal -- in which a backslash is literal -- is skipped as data. So text that only mentions the instruction or the callee -- in a comment, a data or metadata string, a `.file` path, or another symbol's quoted name -- no longer counts, and the quoted, escaped, `unwind`, spaced and comment-separated spellings of a real one do. Thirty-one controls across four tests, each a line from a module `llvm-as` from the pinned toolchain accepts. All but one also lower through `llc`; the exception is a global whose quoted name contains a space and `@`, which is valid IR that NVPTX cannot emit, and which tests the lexer rather than lowering. The PTX controls are also checked with ptxas, where the float cases assemble at sm_100a and are refused at sm_80 as `redux.f32`. Twelve mutations of the lexers each fail a test. The float intrinsic spelling was never matched by the old predicate, but that is not a live mis-target: those are catalog intrinsics, and `select_target_with_generated` takes sm_100a from their recorded requirements regardless -- `redux_f32` selects sm_100a with nothing pinned under either version. Signed-off-by: Vyncint Ng <115854244+vyncint@users.noreply.github.com>

…reduce Device code had no way to ask what target it was being compiled for, so `warp_reduce` could not take the one-instruction `redux.sync` form even when the build was certain to be Ampere or newer: a full warp always paid five shuffles plus five combines. Measured on an A10G over 1M threads reducing 64 times each, that is 159.4 us per launch against 33.6 us, 4.74x, for bit-identical results. A cfg is the mechanism that works here. It is decided before rustc runs, so rustc drops the branch that does not apply and the requirement scan never sees an instruction the target cannot run -- unlike a `__CUDA_ARCH__`-style intrinsic folded at lowering, where the scan reads `redux.sync` out of the dead branch and raises the module floor before anything folds. `cuda-device` gains a build script that resolves the floor and emits `cuda_oxide_sm_at_least = "XY"` for every rung at or below it. `cfg` compares strings rather than numbers, so the rungs are cumulative and device code tests the one its instruction needs. The floor rule lives in `cuda-target-spec`, shared with the backend rather than restated, and follows the backend's own precedence: `CUDA_OXIDE_TARGET`, else `CUDA_OXIDE_DEVICE_ARCH`, else sm_80. Because both sides read the same two variables through the same function, the cfg floor and the selected target cannot disagree. The hint is read with the `DeviceArch` grammar the backend validates it with; one that does not parse is left out, because the backend rejects it before selecting a target. A named target that does not parse fails the build instead of falling back: guessing high emits instructions the target cannot run, guessing low silently drops the faster form. Rust has no specialization, so the reduction reaches a form that exists for only some `(Op, T)` pairs through a provided `ReduceOp::reduce_full_warp` hook, `None` by default. Nine integer pairs override it under the Ampere rung; `f32` keeps the butterfly, since there is no `f32` `redux` before sm_100. Only `N == 32` consults it, leaving the sub-warp path and its out-of-tile lane substitution untouched. One behaviour change: with no explicit target the floor is sm_80, so an integer `warp_reduce` now emits `redux.sync` and its module requires Ampere. Building with `--arch sm_75` keeps the portable form. The `warp_reduce` rustdoc and the warp-level-programming chapter both said this selection did not happen and pointed at this issue; both now describe what it does. Adding the build-dependency puts the first `cuda-device` -> `cuda-target-spec` edge in the graph, and `cuda-device` is in every kernel crate's. So 234 example locks plus the `cuda-macros` device-only test lock each gain that one line; `check-example-license-policy.sh` and `check-device-only-build.sh` fail on a stale lock, so the churn is mandatory rather than cosmetic. It is exactly one added line per file and no version moves. A build-dependency reaches only the host side of the build, so no kernel graph changes. The floor is a lower bound on target selection, and the backend now enforces it. It did not before: `resolve_sm_floor` was consulted only by the build script, so with nothing pinned the ladder advertised sm_80 while a module whose features need only sm_75 -- `ldmatrix` alone -- was built for sm_75, below the capability device code had been compiled against. Both resolvers, native PTX and NVVM, now raise such a selection to the floor, which every architecture at or above it can host. A floor that cannot host the module at all is reported rather than resolved to something else; no path reaches that today, and it is there so a future floor cannot select an architecture family silently. The hint rule is the other half, and it is a correction to the design sketched on the issue. A detected device cannot raise the floor. `CUDA_OXIDE_DEVICE_ARCH` is advisory by contract -- the backend "builds for the detected GPU only when that GPU can actually run the kernel", and otherwise for the arch the kernel requires -- so it says nothing about how high selection will go: a detected sm_120 with a WGMMA module lands on sm_90a, and a floor of 120 taken from it would be a bound the backend is documented to break. What the hint does bound is how low selection may go, since a device that can run the kernel is selected as-is. So it lowers the floor and never raises it: running on Turing still produces sm_75 rather than PTX that device cannot load, and device code is told it has only the rungs that PTX will have. One behaviour change falls out, stated because it reaches builds that never use this cfg: with nothing pinned, a module whose features ask only for sm_75 is now built for sm_80. That is the documented Ampere+ minimum and the same default a featureless module already got; `--arch sm_75` still selects sm_75. Signed-off-by: Vyncint Ng <115854244+vyncint@users.noreply.github.com>

…controls `warp_reduce` now takes `redux.sync` for the nine integer pairs at the Ampere floor and the shuffle butterfly below it, but no example demonstrated either: the only caller of the generic API is `coop_groups_demo`, which also uses thread-block clusters and so cannot be built below sm_90 at all. There was nowhere to read the PTX difference, and nothing that would notice if the fast path started answering differently from the tree it replaces. This example uses one warp and nothing else, so it builds at both floors: cargo oxide build warp_reduce_redux --arch sm_80 redux.sync, no shfl cargo oxide build warp_reduce_redux --arch sm_75 shfl butterfly Four kernels, three of them controls. The `u32` kernel takes every supported op, with inputs that pin the whole warp: a distinct bit per lane, so `or` and `xor` are `0xFFFFFFFF` and `and` over the complements is 0 -- a reduction that dropped lanes could not produce those. The `i32` kernel straddles zero, because signedness is what a single-instruction form gets wrong silently: `redux.sync.min.s32` and `.u32` disagree on exactly those inputs. The two controls are the cases that must *not* move. `f32` has no `redux` before sm_100, so a hook that claimed one would name an instruction that does not exist. A `WarpTile<16>` must reduce within its own tile: the fast path is gated on `N == 32`, and if that gate broke, both tiles would answer 496 instead of 120, which the expected values catch. Signed-off-by: Vyncint Ng <115854244+vyncint@users.noreply.github.com>

Signed-off-by: nihalpasham <nihalp@nvidia.com>

Since NVIDIA#1256 the backend validates CUDA_OXIDE_DEVICE_ARCH once, at its input boundary, and hands selection a DeviceArch: a hint that does not parse is an error before either backend chooses a target. The three readings of a malformed hint this stack used to pin -- the floor resolver and the PTX selector ignoring it, the NVVM selector rejecting it -- no longer exist, so that test was dropped rather than carried forward. What the floor depends on now is narrower. cuda-device's build script reads the same variable with `parse::<DeviceArch>().ok()`, so its floor matches selection exactly when the boundary accepts the values that parse and rejects the rest. The new test asks both about the same inputs -- valid, blank, padded, bare, too short, virtual and family spellings -- and fails if they disagree. Making the boundary trim its input fails it on `" sm_86"`, which the build script would leave out while the backend selected it. The resolver's doc no longer says malformed hints are ignored, since the hint now reaches it parsed. Signed-off-by: Vyncint Ng <115854244+vyncint@users.noreply.github.com>

copy-pr-bot · 2026-10-11T06:06:51Z

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

vyncint · 2026-10-11T06:07:15Z

Rebased onto the mono-repo main and onto #1304's rebased ed917c52. This was not a mechanical rebase, because #1256 (merged 2026-09-23) changed how the device hint reaches selection. The stack is now ed917c52 (#1304) ← e5bc97d4 ← e8f6b24f ← 7ed45774 (yours) ← ab109841.

old head (pre-#1256 base): raw hint ─> floor resolver and PTX ignore a malformed one; NVVM rejects it
new main (#1256):          raw hint ─> input boundary ─> DeviceArch ─> both selectors      (malformed: one input error)
this head:                 build script ─ parse::<DeviceArch>().ok() ─> resolve_sm_floor  (malformed: left out; the boundary rejects it)

What changed

Outside the lock files, this PR's own diff differs from cf4c55c3's in five files: build.rs, export.rs (one signature), cuda-target-spec's lib.rs, selection.rs and the example's Cargo.toml. The other nine are line-for-line the same: cooperative_groups.rs, cuda-device's manifest, select.rs, target/mod.rs, the detector and its two test files, the book page and the example's source.

resolve_sm_floor takes the parsed hint, Option<&CudaArch>. The backend passes the device its boundary validated. Both selectors' call sites are unchanged; only sm_floor_for_selection's parameter type follows. The build script reads the variable with parse::<DeviceArch>().ok(), so it uses exactly the hints that can reach selection. A hint that does not parse is left out of the floor, because the boundary rejects it before either backend chooses a target.
cf4c55c3 pinned three readings of a malformed hint that no longer exist, so I dropped it rather than carry it forward. Its replacement, ab109841, asks both readers of CUDA_OXIDE_DEVICE_ARCH (the boundary's DeviceArchHint and the build script's grammar) about the same eleven inputs, and fails if they disagree. Making the boundary trim its input fails it on " sm_86", which the build script would leave out while the backend selected it.
Your docs: clarify the shared SM floor contract replays byte for byte. export.rs conflicted only on context: your new doc line sits on top of the retyped signature. The one sentence fix(codegen): reject invalid device architecture hints on both backends #1256 made stale, "malformed hints are ignored", is corrected in ab109841.
Locks: all 236 that contain cuda-device were regenerated with cargo update -p cuda-device --offline, and each gains exactly the one "cuda-target-spec" edge. The old commit touched 230; all 230 are still here, with the SIMT workspace lock now at cuda-oxide/Cargo.lock. The extra six are four examples added since (cross_crate_merged_symbols, debug-tests, image_convolution, sparse_mma_fp8) and two the old commit missed: cutile_inter_kernel/simt and interop_cubin_identity/device contained cuda-device but never got the edge, and cargo metadata --locked refuses both at cf4c55c3. I regenerated rather than merged because the textual merge was wrong twice: it put the workspace lock's edge into the root host lock, which has no cuda-device, and it carried atomics' edit into image_convolution's lock, leaving atomics stale.
The example moved to cuda-oxide/crates/rustc-codegen-cuda/examples/warp_reduce_redux/ and depends on the git-root cuda-core by path, like its siblings since build: depend on the shared host crates by path at tip of tree #1466. Its lock is redux_sum's with the root package renamed, as before.

Unchanged: the floor rule, the reduction hooks and warp_reduce's branch. #1311's partial-warp rework does not reach this path: block_reduce calls warp_reduce::<_, _, 32> only for a complete physical warp, and sends a partial tail through scan_contiguous_warp_prefix. So the full-mask redux.sync is still only issued where all 32 lanes are live.

Verification on ab109841: just check passes, which covers fmt-check, both clippy passes, test, test-cuda, every guard and cargo-deny run, check-intrinsics and doc-check: 5,523 tests, 0 failed. just check-intrinsics upstream/main and just book pass too. Per crate: cuda-target-spec 16, cuda-device 114, cuda-oxide-codegen 1,280. #1304 alone is 1,276, and the four added here are this PR's floor tests. cuda-device's tests compile the Ampere-floor assertion with no environment, CUDA_OXIDE_DEVICE_ARCH=sm_120a and =compute_75, and the below-floor one with CUDA_OXIDE_TARGET=sm_75 and CUDA_OXIDE_DEVICE_ARCH=sm_75; all pass.

cargo oxide build warp_reduce_redux through LLVM NVPTX:

build	`.target`	`redux.sync`	`shfl.sync.bfly`
no target, no hint	sm_80	9	19
`--arch sm_75`	sm_75	0	64

With CUDA_OXIDE_DEVICE_ARCH=compute_75, the build stops at the boundary with invalid CUDA_OXIDE_DEVICE_ARCH `compute_75`: expected sm_<capability> with an optional `a` suffix and writes no PTX.

With ptxas 13.3, the first PTX assembles at sm_80 and the second at sm_75. The first, retargeted to sm_75, is refused with Feature 'redux' requires .target sm_80 or higher.

Two statements in the description now predate the head, and since it is your text I have left it to you: the malformed-hint regression it describes is the dropped cf4c55c3, and the "Remaining prerequisite" is what #1304's 6b4ac7d0 (now ed917c52) addresses. The B200 runs predate this rebase, and I have not repeated any GPU run for it.

vyncint mentioned this pull request Sep 19, 2026

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

Open

nihalpasham requested changes Sep 23, 2026

View reviewed changes

vyncint mentioned this pull request Sep 23, 2026

fix(codegen): match redux.sync's f32 form on the instruction, not on nearby text #1304

Open

vyncint force-pushed the feat/arch-derived-sm-floor-cfg branch from 58b0d79 to 70ede7a Compare September 23, 2026 07:32

vyncint force-pushed the feat/arch-derived-sm-floor-cfg branch from 70ede7a to 08d95ff Compare September 23, 2026 07:45

nihalpasham requested changes Sep 24, 2026

View reviewed changes

vyncint force-pushed the feat/arch-derived-sm-floor-cfg branch from db5c23f to 96ea0b5 Compare September 24, 2026 13:47

nihalpasham requested changes Sep 24, 2026

View reviewed changes

vyncint force-pushed the feat/arch-derived-sm-floor-cfg branch from 96ea0b5 to cf4c55c Compare September 24, 2026 14:36

0z5a mentioned this pull request Sep 26, 2026

bench(cuda-device): add paired warp reduction measurements #1339

Closed

0z5a mentioned this pull request Sep 26, 2026

fix(mir-lower): unify equal-width shift representations #1338

Open

nihalpasham added the cuda-oxide SIMT programming model: rustc backend, cargo oxide, cuda-device, cuda-host, book label Oct 11, 2026

vyncint and others added 4 commits October 11, 2026 12:39

docs: clarify the shared SM floor contract

7ed4577

Signed-off-by: nihalpasham <nihalp@nvidia.com>

vyncint force-pushed the feat/arch-derived-sm-floor-cfg branch from cf4c55c to ab10984 Compare October 11, 2026 06:06

pedrozaz mentioned this pull request Oct 11, 2026

feat(cuda-device): support u64 reductions and scans in cooperative_groups #1497

Open

Navigation Menu

Repository navigation

feat(cuda-device): derive an SM floor cfg and use redux.sync in warp_reduce - #1305

feat(cuda-device): derive an SM floor cfg and use redux.sync in warp_reduce#1305
vyncint wants to merge 5 commits into
NVIDIA:mainfrom
vyncint:feat/arch-derived-sm-floor-cfg

vyncint commented Sep 19, 2026 •

edited by nihalpasham

Loading

Uh oh!

nihalpasham left a comment

Uh oh!

vyncint commented Sep 23, 2026

Uh oh!

nihalpasham left a comment

Uh oh!

vyncint commented Sep 24, 2026

Uh oh!

nihalpasham left a comment

Uh oh!

vyncint commented Sep 24, 2026

Uh oh!

copy-pr-bot Bot commented Oct 11, 2026

Uh oh!

vyncint commented Oct 11, 2026

Uh oh!

This branch has not been deployed

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

vyncint commented Sep 19, 2026 • edited by nihalpasham Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What this adds

Verification

Remaining prerequisite

Uh oh!

nihalpasham left a comment

Choose a reason for hiding this comment

Uh oh!

vyncint commented Sep 23, 2026

Why a detected device cannot raise the floor

Carried into both selectors

One behaviour change, stated because it reaches builds that never use this cfg

Hardware

Uh oh!

nihalpasham left a comment

Choose a reason for hiding this comment

Uh oh!

vyncint commented Sep 24, 2026

Uh oh!

nihalpasham left a comment

Choose a reason for hiding this comment

Uh oh!

vyncint commented Sep 24, 2026

Uh oh!

copy-pr-bot Bot commented Oct 11, 2026

Uh oh!

vyncint commented Oct 11, 2026

Uh oh!

This branch has not been deployed

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

vyncint commented Sep 19, 2026 •

edited by nihalpasham

Loading