Repository navigation
Conversation
nihalpasham
added
the
cuda-oxide
vyncint
added a commit
to vyncint/cuda-oxide
that referenced
this pull request
Oct 30, 2026
…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 257 example and model locks, the SIMT workspace lock and two test-fixture locks 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. No version moves. Three of the example locks (`inline_intent` and the two `error_union_shared_pointer_*`) also gain the `cuda-host` -> `cuda-device` edge they lack on main, because cargo rewrites the whole resolution; NVIDIA#1667 adds the same line. 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>
nihalpasham
added
bug
Collaborator
|
Thanks for catching these three stale example lockfiles. Closing this PR because commit
|
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.
Summary
The committed
Cargo.lockofinline_intent,error_union_shated_pointer_addranderror_union_shared_pointer_layoutdoes not record thatcuda-hostdepends oncuda-device. Cargo rewrites the lockfile when these examples are built, so runnning the repository smoketest leaves the working tree dirty. Same class as #993 / #994Changes
"cuda_device"tocuda-host's locked dependencies in the three lockfiles aboveTesting
cargo metadata --locked --offlinein each example directory fails with "cannot update the lock file ... because --locked was passed".bash cuda-oxide/scripts/sync-example-locks.sh --checkpasses.cuda-oxide/scripts/smoketest.sh -o '^(inline_intent|error_union_shared_pointer_addr|error_union_shared_pointer_layout)$'passes 3/3 (the twoerror_union_*fixtures fail to compile as expected) and leaves the tree clean.--locked --offline; these three are the only stale ones.cutile_inter_kernelandmathdx_ffi_testcould not be checked offline.