Repository navigation
chore: move bloomery to NVIDIA/cuda-rust 9df9e7ce and speed up the CuTe verifier - #4
Merged
midagedev merged 71 commits intoOct 26, 2026
Conversation
Both component trees carried their own CONTRIBUTING.md and LICENSE next to the root ones. This leaves one of each. `cutile-rs/CONTRIBUTING.md` repeated the DCO text and sign-off instructions from the root guide and said nothing about the monorepo. The two things it had that the root guide lacked move there: the branch and PR naming convention (with `docs` rather than `doc`, which is what the history on `main` uses) and how to run the Tile test scripts from `cutile-rs/`. The root guide's Toolchains section, AGENTS.md, and the cutile-rs README now point at the root guide. The LICENSE copies under `cutile-rs/` and `cuda-oxide/` duplicated the root file, and nothing consumed them: every crate declares its license through the SPDX `license` field rather than `license-file`, GitHub reads the root LICENSE, and Apache-2.0 asks for one copy per distribution. The root README's per-component license table collapses to one sentence, and the two cuda-oxide references (its README and THIRD_PARTY_NOTICES) point at the root file. The cuda-oxide copy was added deliberately in the recent scoping commit (bb8d4fa), so this is the place to say so if a local copy is wanted after all. No code changes. `check-source-references.sh` and `check-spdx-headers.sh` pass. Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com> Co-authored-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
…VIDIA#1458) Puts the book badge back in `cutile-rs/README.md`, next to the crates.io badge. NVIDIA#1394 moved both badges to the repository root, but the crates.io page renders this README and a visitor there wants the book link more than anything else on it. The build badge stays at the root with the other CI badges (NVIDIA#1395), which matches how package READMEs in other monorepos carry package-level badges while the root carries CI status. Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com> Co-authored-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
…oot CI badges (NVIDIA#1395) Companion to NVIDIA#1394: the references outside `cutile-rs/` that still pointed at the old NVlabs/cutile-rs home. - `repository` field of cuda-bindings, cuda-core, cuda-core-derive, and cuda-async (crates.io and docs.rs read it). - cuda-oxide glossary, VMM page, cuda-host README, scaffold template, and doctor comments now say the shared host crates live at the repository root; the four rows in `dependency-licenses.csv` record them as 0.4.0 path dependencies instead of 0.3.1 from crates.io. - Root README gets CI badges, one per pipeline, plus a docs badge. A single combined badge is proposed separately in NVIDIA#1457. The cuda-oxide book builds warning-free, and the source-reference, host-API-path, and pin-parity guards pass. --------- Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com> Co-authored-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
## Summary A later push to `main` waits for the previous run of a workflow instead of cancelling it. Other refs still cancel the older run. ## Changes - Set `cancel-in-progress` from the ref name so `main` runs finish. This covers CI, docs, examples, status-guard, naming-guard, CodeQL, and the Tile workflow. The book workflow already left `main` runs in place. ## Testing Workflow configuration only. No crate, example, or guard behavior changed, so `just check` and `cargo oxide run` were not run. - [ ] `just -f cuda-oxide/Justfile check` passes (the local mirror of CI: fmt, clippy, tests, guards, docs) - [ ] `cargo oxide run <example>` passes, or `cuda-oxide/scripts/smoketest.sh -o '^<example>$'` - [ ] New example added (if applicable) ## Checklist - [x] All commits signed off (`git commit -s`) - [x] SPDX headers on new source files --- > Questions about the review? Ping us in [#contributors on Discord](https://discord.gg/ZUEr4AhH5C). Signed-off-by: Roman Ivanov <rivanov@nvidia.com>
## Summary Skip Tile, CodeQL, and the naming guard when a change cannot affect what those jobs read. ## Changes - Declare the Tile `any_changed` filter so a non-Tile diff skips the Tile jobs, while a push with no detectable diff still runs them. The job condition already read `any_changed`, but `dorny/paths-filter` only sets an output for a filter that is declared. That name was not declared, so the output stayed empty. Empty is not `'true'`, and the condition was `any_changed != 'true'`, so it was always true and every Tile job ran on every push. - Run CodeQL on push and pull request only when Rust, Cargo manifests, `.github/`, shell, or Python files change. The weekly schedule still scans the whole tree. - Run the naming guard only for `cuda-oxide/crates/`, its script, and its workflow. The comment records why `cutile-rs` and the shared host crates stay out: the script's `SEARCH_ROOTS` does not read them. ## Testing Workflow configuration only. No crate, example, or guard behavior changed, so `just check` and `cargo oxide run` were not run. - [ ] `just -f cuda-oxide/Justfile check` passes (the local mirror of CI: fmt, clippy, tests, guards, docs) - [ ] `cargo oxide run <example>` passes, or `cuda-oxide/scripts/smoketest.sh -o '^<example>$'` - [ ] New example added (if applicable) ## Checklist - [x] All commits signed off (`git commit -s`) - [x] SPDX headers on new source files --- > Questions about the review? Ping us in [#contributors on Discord](https://discord.gg/ZUEr4AhH5C). Signed-off-by: Roman Ivanov <rivanov@nvidia.com>
…IDIA#1457) Stacked on NVIDIA#1395; only the last commit is new, and the diff shrinks to it once NVIDIA#1395 merges. Badges are per workflow file, and the SIMT and Tile pipelines have to stay separate (the Tile lanes run on self-hosted runners and cannot trigger on pull_request), so the root README ends up with two CI badges. `ci-status.yml` runs whenever `CI` or `cutile-rs-pr` completes on `main`, takes the newest completed push run of each pipeline that actually executed jobs, and fails if either failed. The README then carries one CI badge pointing at it. Runs whose jobs were all skipped (the Tile `changes` gate on a non-Tile push) and cancelled runs are skipped over, so a later no-op run cannot turn the badge green while the last real run was red. The logic was dry-run against this repo's run history. The badge reads "no status" until the first run after merge. The badge row is centered, matching the component READMEs. --------- Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com> Co-authored-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
…A#1466) First PR under NVIDIA#1459. cuda-oxide depends on the shared host crates by path, but the dependency lines also carried `version = "=0.4.0"`, and `check-shared-crate-pin.sh` required that exact string in all 234 example workspaces and in the `cargo oxide new` scaffold. Every host-crate release therefore meant editing cuda-oxide's manifest, every example manifest, the scaffold constant, and the license inventory before the tree was consistent again. This drops the version requirement from every in-tree path dependency on cuda-bindings, cuda-core, and cuda-async: the three workspace lines and 238 lines across the example manifests. A path dependency with no version always resolves to the sibling directory, so cuda-oxide now tracks the host crates at tip of tree. Nothing in this tree that depends on them (cuda-host, cuda-macros, the examples) is published, so nothing needs the version. Lockfiles are unchanged; the touched example workspaces still pass `cargo metadata --locked`. `check-shared-crate-pin.sh` keeps its name and its CI job but checks the new invariant: the four root crates agree on a version, in-tree manifests name the host crates as path dependencies with no version, and the scaffold's `SHARED_HOST_CRATES_VERSION`, the one place a crates.io version lives (out-of-tree projects cannot use the git-root paths), equals the root crates' version. It rejects a reintroduced version field; I checked that by hand. Deriving the scaffold constant from the root manifests at build time is the next PR under NVIDIA#1459. The four host-crate rows in `dependency-licenses.csv` now record `path` instead of a version number; `check-dependency-licenses.sh` compares names only, so they remain valid. Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com> Co-authored-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Update CUDAoxide from `nightly-2026-08-28` to `nightly-2026-10-11`. This brings in rust-lang/rust#163242, needed for the separate kernel parameter-attribute work in NVIDIA#1483 (related to NVIDIA#1088 and NVIDIA#1098). ## Changes - Adapt to rustc API changes, generic floating-point intrinsics, and integer min/max. - Follow the new `core::intrinsics::reflection::type_id` path in the kernel-name helper. - Update toolchain/Nix pins, diagnostic snapshots, and intrinsic evidence. Intrinsic declarations are unchanged. ## Validation - CI: 55 passed, 7 expected skips; all 234 examples compile. - Compiler tests, formatting, clippy, and IR comparison pass. Targeted checks pass on both RTX 5090s with no memcheck errors. - GPU sweep: 232/233 passed. The existing FP8 example failure is addressed separately in NVIDIA#1487; hardware-limited cases retain compile-only coverage. The existing scaffold issue—mixing registry and Git runtime types—also remains a separate follow-up. ## Checklist - [x] All commits signed off - [x] SPDX headers on new source files, where applicable --------- Signed-off-by: nihalp <nihalp@nvidia.com>
Fix `sparse_mma_fp8`, which tested twelve metadata encodings even though plain FP8 `mma.sp` permits only six. The invalid encodings caused 5,104 mismatches on RTX 5090. The example and layout probe now use the six valid encodings. Inline comments and generated safety docs explain the restriction and link to the [PTX specification](https://docs.nvidia.com/cuda/archive/12.4.1/parallel-thread-execution/index.html#warp-level-sparse-matrix-storage). Instruction lowering and ABI are unchanged. Also set `rust-analyzer.cargo.features` to `[]` so VS Code stops enabling every optional feature. ## Validation - Both regressions fail before the fix; 319 generator tests and 3 example tests pass. - Both RTX 5090s pass all 6,144 results, with zero memory or synchronization errors. - Verified with both the old nightly and `nightly-2026-10-11` after NVIDIA#1482; generated outputs and all 1,029 ABI entries pass. - Before the editor-only update: CI had 54 passes and 6 expected skips; formatting and strict clippy passed. The updated settings file validates; new CI results are pending. Follow-up to NVIDIA#1295. Running examples on GPUs in CI remains separate work. --------- Signed-off-by: nihalp <nihalp@nvidia.com>
Use rustc's `noalias` and proven `dereferenceable(N)` facts on `cuda-oxide` kernel references. This lets LLVM reuse input reads across writes to separate outputs. Raw pointers and `DisjointSlice` storage receive no new attributes. Rust's interior-mutability, pinning, and lifetime rules still apply. Explicit `readonly` is deferred because current NVPTX backends can lose volatile reads. Includes regression tests, corrected atomic example pointer contracts, and load-width checks that accept cached loads. Validation: compiler/backend suites and strict checks pass; reference launches pass memory, race, and synchronization checks on both RTX 5090s. LLVM and both NVVM paths preserve two volatile reads while combining ordinary reads. Final hosted CI is pending. The existing FP8 failure is handled by NVIDIA#1487. The existing AtomicF16 sanitizer failure remains a separate follow-up. Partially addresses NVIDIA#1088 and NVIDIA#1098. ## Checklist - [x] All commits signed off - [x] SPDX headers on new source files --------- Signed-off-by: nihalp <nihalp@nvidia.com>
## Summary
Successive pushes to `main` share one concurrency group per workflow, so
the newer run waits until the older one finishes. Each commit on `main`
now has its own group and starts immediately. The previous run still
finishes. Any other ref still cancels its older run.
The book workflow is the exception. Its builds stay in one group per ref
so a slower older commit cannot publish after a newer one and replace
the site.
## Changes
- Append the commit SHA to the concurrency group when the ref is `main`,
in CI, CodeQL, docs, examples, status-guard, naming-guard, and the Tile
workflow.
- Leave `book.yml` as one group per ref. The deploy job keeps the shared
`pages` group.
- Keep `cancel-in-progress: ${{ github.ref_name != 'main' }}`. The SHA
in the group already separates commits on `main`, so this line is what
states the policy in plain sight: runs on `main` are not cancelled, and
every other ref still cancels the older run. Removing it would leave
that rule buried in the group string.
## Testing
Workflow configuration only. No crate, example, or guard behavior
changed, so `just check` and `cargo oxide run` were not run.
- [ ] `just -f cuda-oxide/Justfile check` passes (the local mirror of
CI: fmt, clippy, tests, guards, docs)
- [ ] `cargo oxide run <example>` passes, or
`cuda-oxide/scripts/smoketest.sh -o '^<example>$'`
- [ ] New example added (if applicable)
## Checklist
- [x] All commits signed off (`git commit -s`)
- [x] SPDX headers on new source files
---
> Questions about the review? Ping us in [#contributors on
Discord](https://discord.gg/ZUEr4AhH5C).
---------
Signed-off-by: Roman Ivanov <rivanov@nvidia.com>
…VIDIA#1499) Adds the AI-assisted contributions section from NVIDIA's open source CONTRIBUTING template, verbatim, placed before the IP review section as in the template and listed in the table of contents. Documentation only. Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com> Co-authored-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
…NVIDIA/cuda-rust (NVIDIA#1490) ## What this adds The SIMT tree still named the pre-rename repository in its Cargo metadata, its license inventory and the device-codegen ICE message, while the four host crates already pointed at `NVIDIA/cuda-rust` (NVIDIA#1395). This change moves the SIMT tree to the same home: one inherited `[workspace.package] repository` line for the 26 workspace members, explicit fields on the two first-party crates outside that workspace, 29 inventory rows and the ICE text. The only compiled change is the string in that diagnostic. ```text cargo metadata `repository`, 28 SIMT crates before: https://github.com/NVlabs/cuda-oxide.git after: https://github.com/NVIDIA/cuda-rust (one value) first-party inventory rows at the old URL before: 29 of 33 after: 0 of 33 ICE text before: ...Please file at https://github.com/NVlabs/cuda-oxide/issues after: ...Please file at https://github.com/NVIDIA/cuda-rust/issues ``` ```rust // The only compiled change: the diagnostic printed when device codegen panics. tcx.dcx().fatal(format!( "[rustc_codegen_cuda] Internal compiler error in device codegen: {msg}. \ This is a bug in cuda-oxide. Please file at https://github.com/NVIDIA/cuda-rust/issues" )); ``` ## What we fixed - `cuda-oxide/Cargo.toml` `[workspace.package] repository`, inherited by every `repository.workspace = true` member (26 crates). - `rustc-codegen-cuda/Cargo.toml` and `examples/mathdx_ffi_test/libmathdx-sys/Cargo.toml`, the two first-party manifests that cannot inherit the workspace field. - 29 of the 33 first-party rows of `cuda-oxide/dependency-licenses.csv`; the four git-root host rows already carried the new URL. - The fatal internal-compiler-error message in `rustc_codegen_cuda` now names `https://github.com/NVIDIA/cuda-rust/issues`. <details><summary>Implementation</summary> - The URL form matches the host crates (`https://github.com/NVIDIA/cuda-rust`, no `.git` suffix). - `repository` is not part of `Cargo.lock` and nothing in the tree reads it, so no lockfile, symbol or PTX changes. - The inventory guard compares crate names only; the rows are hand-maintained, so the URL column was edited directly. - The remaining old-URL sites under `cuda-oxide/` are the cargo-oxide clone and scaffold constants (NVIDIA#1489), the book and crate README links (NVIDIA#1491), and issue and pull permalinks that keep their historical form. </details> ## Verification Checked at `dad294fd` (merged with `main` at `bc46b1ff`): - ✅ `git grep NVlabs/cuda-oxide` over the three manifests, the inventory and `lib.rs`: 1, 1, 1, 1, 29 hits on `main`, 0 after the change. - ✅ `cargo metadata --locked --offline --no-deps` in the SIMT workspace (26 members), the backend crate workspace and the `mathdx_ffi_test` workspace: one `repository` value, `https://github.com/NVIDIA/cuda-rust`; `git status` clean afterwards. - ✅ `strings` on the rebuilt `librustc_codegen_cuda.so`: `Please file at https://github.com/NVIDIA/cuda-rust/issues` once, the old URL zero times; `main`'s backend still shows the old URL. - ✅ `cargo oxide run vecadd` and `cargo oxide run mathdx_ffi_test --emit-nvvm-ir --arch=sm_120` on an RTX 5090 through the rebuilt backend: `SUCCESS` and `Passed: 9 / Failed: 0`; compute-sanitizer memcheck 0 errors on both; `vecadd` PTX identical to `main` modulo symbol hashes. - ✅ All 18 `scripts/check-*.sh` guards, `sync-example-locks.sh --check`, `cargo deny check` for cargo-oxide, rustc-codegen-cuda, the device-only fixture and the root workspace, `cargo oxide fmt --check`, clippy with `-D warnings`, and the unit, doc and rustdoc suites (5479 tests passed, 0 failed). - ✅ No file in this PR appears in NVIDIA#1489 or NVIDIA#1491. ## Caveats - The `oxide-artifacts` inventory row now names the current home; the published 0.2.1 crate still declares the old URL until the next publish from `crates/oxide-artifacts`. - No test reaches the ICE message; the string change is confirmed by compilation and the dylib strings only. - Clone and scaffold URLs in cargo-oxide and the book and README links are left to NVIDIA#1489 and NVIDIA#1491. Signed-off-by: Alejandro Soto Franco <sotofranco.eng@gmail.com>
…NVIDIA#1491) ## What this adds The SIMT book, its `conf.py` and the `cargo-oxide`, `ptx-parse`, `ptx-schedule` and `pinned_overlap` READMEs still named `NVlabs/cuda-oxide` in 23 current-repository references and reached the repository only through GitHub's redirect. This moves every one of them to `NVIDIA/cuda-rust`: four `cargo install --git` lines, the clone step, 16 tree and blob links to 13 targets, the Discussions link and the book's GitHub icon. Issue and pull-request permalinks keep their historical form, as AGENTS.md asks. ```text Before: git clone https://github.com/NVlabs/cuda-oxide.git # redirect; creates cuda-oxide/ holding the monorepo cd cuda-oxide After: git clone https://github.com/NVIDIA/cuda-rust.git cd cuda-rust cd cuda-oxide # build, doctor and run live here ``` ```bash cargo +nightly-2026-10-11 install --git https://github.com/NVIDIA/cuda-rust.git cargo-oxide ``` ## What we fixed - `conf.py`: the GitHub icon links `https://github.com/NVIDIA/cuda-rust`. - Four `cargo install --git` lines (cargo-oxide README, Building from Source, Writing Your First Kernel, Installation) and the clone step in Building from Source. - 16 tree and blob links to 13 targets across seven book pages and the `ptx-parse`, `ptx-schedule` and `pinned_overlap` READMEs. - The Discussions link in the ecosystem appendix. - Building from Source: the build, verify and common-commands sections run from `cuda-oxide/`, where the nightly pin and the `cargo oxide` alias apply, and the book and rustdoc sections start from the repository root. The shared host runtime is described as a path dependency on the repository-root crates, documented from the root workspace. - `pinned_overlap` README: the run line names the cuda-oxide root, like the other example READMEs. <details><summary>Implementation</summary> - Documentation sources only: 11 book pages, `conf.py` and four READMEs. No code, manifest or lockfile changes. - The tree and blob links already carried monorepo paths with the `cuda-oxide/` prefix; only the repository name changed. - Issue and PR permalinks, 18 lines in 15 files under `cuda-oxide/`, are untouched. - Code constants, the three `repository` manifest fields, the dependency-license table and the ICE message's issue URL move in NVIDIA#1489 and NVIDIA#1490. </details> ## Verification Checked at `e678f228`: - ✅ All 13 distinct link targets exist on `origin/main` (`git ls-tree`), tree links to directories and blob links to files. - ✅ `NVIDIA/cuda-rust` has Discussions enabled; `NVlabs/cuda-oxide` still redirects there. - ✅ `sphinx-build -W --keep-going -b html` (Sphinx 9.1.0, 41 pages) builds with no warnings; the rendered HTML has no non-permalink `NVlabs/cuda-oxide` URL left (24 in the book and READMEs on main before, 1 after, owned by NVIDIA#1489). - ✅ All 18 `scripts/check-*.sh` guards pass, including `check-source-references.sh`, `check-book-api-names.sh` and `check-host-api-paths.sh`. - ✅ `cargo test` for cargo-oxide, ptx-parse and ptx-schedule (243, 35 and 27 tests), clippy with `-D warnings`, rustdoc, `cargo oxide fmt --check` and `sync-example-locks.sh --check` pass. - ✅ `cargo oxide run pinned_overlap` on an RTX 5090 (sm_120a) prints its success marker; compute-sanitizer memcheck reports 0 errors. - ✅ A dry-run `cargo install --git` against a snapshot of this tree discovers `cargo-oxide` in the nested workspace. - ✅ From the monorepo root, `cargo oxide doctor` reports the stable 1.98.0 toolchain and `cargo oxide build vecadd` fails on nightly-only flags; from `cuda-oxide/` both succeed. The directory anchors were re-checked with the same book build and guards. ## Caveats - The 18 kept permalinks rely on GitHub's redirect, by policy. - No guard checks repository URLs; the link targets were verified by hand against `origin/main`. - The remaining `NVlabs/cuda-oxide` references live in cargo-oxide code constants, the three `repository` manifest fields, the dependency-license table and the ICE message, owned by NVIDIA#1489 and NVIDIA#1490. --------- Signed-off-by: Alejandro Soto Franco <sotofranco.eng@gmail.com> Signed-off-by: nihalpasham <nihalp@nvidia.com> Co-authored-by: nihalpasham <nihalp@nvidia.com>
…t source (NVIDIA#1488) ## What this adds `cargo oxide new` wrote a manifest that took `cuda-device` and `cuda-host` from the repository's git URL but `cuda-core` (and `cuda-async` with `--async`) from crates.io 0.4.0. `cuda-host` depends on `cuda-core` by path inside that checkout, so every scaffolded project resolved two `cuda-core` packages and the code `#[cuda_module]` generates failed to type-check. The templates now name every CUDA crate from the same git source, so cargo resolves one copy of each. ```text Before: cuda-host = { git = <repo> } cuda-core = "0.4.0" Cargo.lock: cuda-core x2 (registry + git), cuda-bindings x2, cuda-core-derive x2 error[E0308]: expected `cuda_core::simt::context::CudaContext`, found `CudaContext` 7 errors sync, 13 errors async; nothing runs After: cuda-host = { git = <repo> } cuda-core = { git = <repo> } cuda-async = { git = <repo> } Cargo.lock: cuda-core x1, cuda-bindings x1, cuda-core-derive x1, all at one commit PASSED: all 1024 elements correct ``` ```rust #[cuda_module] mod kernels { #[kernel] pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) { /* ... */ } } let module = unsafe { kernels::load(&ctx)? }; // generated code names cuda-host's CudaContext ``` ## What we fixed - Both templates take `cuda-core` and `cuda-async` from `GIT_REPO`, the source `cuda-host` comes from; `SHARED_HOST_CRATES_VERSION` is gone. - `check-shared-crate-pin.sh` section 3 requires `git = "{GIT_REPO}"` on every host-crate line in the templates and rejects a crates.io requirement there. - New test `scaffold_takes_every_cuda_crate_from_one_git_source` asserts every `cuda-*` line in both templates names the same git source as `cuda-device`. - The manifest comment and the status-guard step name that still described a crates.io scaffold now describe the git source. ## Scaffold dependency contract Every CUDA crate in a new project comes from one git source with no version requirement. The first build locks all of them to one commit of that repository, and `cargo update` moves them together. A crates.io `cuda-core` or `cuda-async` beside a git `cuda-host` is always a second crate and does not compile. <details><summary>Implementation</summary> - The regression began at bd6c3ce (NVIDIA#29, 2026-10-01), which moved `cuda-oxide` from crates.io `cuda-core` 0.3.1 onto the git-root path copies while the scaffold kept its crates.io constant; NVIDIA#1466 only dropped the `version = "=0.4.0"` qualifier. Lockfiles resolved at the commit before NVIDIA#29 hold one `cuda-core`; at NVIDIA#29, before NVIDIA#1466 and at NVIDIA#1466 they hold two. - Cargo identifies a package by name, version and source. A path dependency declared inside a git-sourced crate resolves to that git source, and nothing unifies it with a registry copy, so the scaffold has to name `cuda-core` from the source `cuda-host` uses. - This is the pre-NVIDIA#1102 form of the template minus `cuda-bindings`, which no user code names and which `cuda-core` pulls by path. - The failures are identity mismatches on `CudaContext`, `CudaModule`, `DeviceBuffer` and the `DriverError` closure: 6 E0308 plus 1 E0277 in the sync template; 9 E0308, 3 E0277 and 1 E0599 in the async one. - This supersedes the NVIDIA#1466 plan to derive the scaffold's version from the root manifests: with no version in the template there is nothing to derive. </details> ## Verification Checked at `180770a9`: - ✅ `cargo test -p cargo-oxide --all-targets`: 244 passed. With the old template lines restored and the new tests kept, 3 scaffold tests fail. - ✅ `cargo clippy -p cargo-oxide --all-targets -- -D warnings`, `cargo fmt --check` and rustdoc with `-D warnings` clean. - ✅ All 18 `check-*.sh` guards pass. `check-shared-crate-pin.sh` rejects four reintroduced forms (`cuda-core = "0.4.0"`, `{ version = "0.4.0" }`, a different git URL, `cuda-async = "0.4.0"`) with exit 1. - ✅ Baseline: main's `cargo oxide new` for both templates fails to build, 7 and 13 errors, with two `cuda-core`, `cuda-bindings` and `cuda-core-derive` per lockfile (and two `cuda-async`). - ✅ Fix: `cargo oxide new` then `cargo oxide run` for both templates on an RTX 5090 (sm_120a, CUDA 13.4): `PASSED: all 1024 elements correct`, `compute-sanitizer --tool memcheck` 0 errors, one `cuda-core` per lockfile and zero registry `cuda-*` packages. Same result on the author's RTX 5060 Laptop (sm_120, CUDA 13.3). - ✅ Merges cleanly with NVIDIA#1489 in either order; after both, every template line names `https://github.com/NVIDIA/cuda-rust.git`. ## Caveats - The test and guard compare the git URL only. A `rev`, `branch` or `tag` added to one line still resolves a second `cuda-core`, even at the same commit; a follow-up tightens the comparison to the full source. - No CI job compiles a scaffolded project, which is how the broken template survived from 2026-10-01; a scaffold compile gate is a separate follow-up. - The GPU runs used a cached backend; the standalone backend build from the scaffold's git checkout is unchanged by this PR and was not re-run here. - `GIT_REPO` still names `NVlabs/cuda-oxide`, which GitHub redirects; NVIDIA#1489 renames it. Closes NVIDIA#1459. --------- Signed-off-by: Alejandro Soto Franco <sotofranco.eng@gmail.com> Signed-off-by: nihalpasham <nihalp@nvidia.com> Co-authored-by: nihalpasham <nihalp@nvidia.com>
## What this adds
`cargo oxide new` writes the repository address of `cuda-device` and
`cuda-host` into every new project's manifest, and the backend fallback
clone fetches the same address when no cuda-oxide dependency pins a
commit. Both still named the pre-move `NVlabs/cuda-oxide` repository,
which resolves only through GitHub's redirect and leaves every new
lockfile recording the old source. Both now name `NVIDIA/cuda-rust`.
```text
cargo oxide new demo
Before cuda-device = { git = "https://github.com/NVlabs/cuda-oxide.git" }
source = "git+https://github.com/NVlabs/cuda-oxide.git#bc46b1ff..." (via redirect)
checkout ~/.cargo/git/checkouts/cuda-oxide-<hash>/bc46b1f/
After cuda-device = { git = "https://github.com/NVIDIA/cuda-rust.git" }
source = "git+https://github.com/NVIDIA/cuda-rust.git#bc46b1ff..." (direct)
checkout ~/.cargo/git/checkouts/cuda-rust-<hash>/bc46b1f/
```
```bash
cargo oxide new demo && cd demo
cargo oxide build # builds the backend from the checkout Cargo made for cuda-device
```
## What we fixed
- `GIT_REPO` in `scaffold.rs`, interpolated into both scaffold
templates, now reads `https://github.com/NVIDIA/cuda-rust.git`.
- The fallback `git clone --depth 1` in `backend.rs`, used only for
projects with no cuda-oxide dependency, clones the same address.
- The `backend_source.rs` module diagram, its test constant and the
matching README diagram follow: the checkout is `cuda-rust-<hash>`, and
the backend crate and the toolchain pin sit under `cuda-oxide/` in it.
- A scaffold test asserts the emitted `git =
"https://github.com/NVIDIA/cuda-rust.git"` lines in both templates and
rejects the old address.
<details><summary>Implementation</summary>
- Backend discovery never reads the URL:
`dependency_source_from_metadata` keys on the package name and
`manifest_path`, `git_source_rev` keeps only the `#<sha>` fragment, and
`checkout_root` walks up to the first directory holding `Cargo.toml`
plus the codegen crate, bounded by `.git` or `.cargo-ok`.
- `codegen_crate_in_checkout` prefers
`cuda-oxide/crates/rustc-codegen-cuda` and falls back to the historical
flat `crates/rustc-codegen-cuda`; `pinned_channel` walks from the
codegen crate upward, so the component nightly wins over the stable pin
at the git root. The same two functions serve the fallback clone.
- The backend cache is keyed by commit (`source-rev.txt`), not by URL,
so cached backends and existing projects are unaffected.
- The historical issue permalink to `NVIDIA/issues/96` in the
README stays, per the repository guide.
</details>
## Verification
Checked at `746d60eb`:
- ✅ `cargo test -p cargo-oxide --all-targets`: 243 passed; 244 with the
new URL test.
- ✅ `cargo clippy -p cargo-oxide --all-targets -- -D warnings`, `cargo
fmt --check`, rustdoc with `-D warnings`, and all 18
`scripts/check-*.sh` guards pass.
- ✅ `grep -rn NVlabs crates/cargo-oxide/src` returns nothing; the
remaining hits are the README install line (NVIDIA#1491) and the NVIDIA#96
permalink.
- ✅ In an isolated `CARGO_HOME`: `cargo oxide new` writes the new
address in both templates; `cargo generate-lockfile` records
`git+https://github.com/NVIDIA/cuda-rust.git#bc46b1ff` for
`cuda-device`, `cuda-host` and `cuda-macros` in a `cuda-rust-*`
checkout; `cargo oxide build` builds the backend from that checkout's
`cuda-oxide/crates/rustc-codegen-cuda` under nightly-2026-10-11 and
caches it with `source-rev.txt = bc46b1f`.
- ✅ The old-address baseline binary scaffolds `NVlabs/cuda-oxide.git`
and resolves the same commit through the redirect; `git ls-remote` on
both URLs returns the same HEAD.
- ✅ A shallow clone of `NVIDIA/cuda-rust` builds the backend offline
from its nested codegen crate under the component nightly, which is the
fallback route.
- ✅ Auto-merges with NVIDIA#1488, NVIDIA#1490 and NVIDIA#1491 in either order.
## Caveats
- A scaffolded project still fails to compile at this head with two
`cuda-core` crates in its graph; that is NVIDIA#1488's fix and is independent
of the address.
- A graph that mixes the old and new address for the same crate resolves
two checkouts and `cargo oxide` refuses it;
`[patch."https://github.com/NVlabs/cuda-oxide.git"]` pointing the crate
at the new address collapses the graph to one source.
- Machines that already hold the old-address git database fetch the
repository once more under the new name; the backend cache is reused
because it is keyed by commit.
- The remaining old-address mentions are the install lines (NVIDIA#1491) and
the manifest `repository` fields and license CSV (NVIDIA#1490).
---------
Signed-off-by: Alejandro Soto Franco <sotofranco.eng@gmail.com>
Signed-off-by: nihalpasham <nihalp@nvidia.com>
Co-authored-by: nihalpasham <nihalp@nvidia.com>
…tion (NVIDIA#1454) ## What this adds A tuple holding a cluster-shared pointer from `cluster::map_shared_rank` now compiles. Tuple construction casts each element to the tuple's declared element type before building the aggregate, the same normalization that struct, enum, array, union and slice construction already apply, so the `addrspace(7)` pointer enters the generic slot through an `addrspacecast` instead of failing MIR verification. ```text Before: Verification failed for 'test_dsmem_ring_exchange': MirConstructTupleOp operand 0 type mismatch. Expected: mir.ptr <ui32, addrspace:0, kind:RawConst> Actual: mir.ptr <ui32, addrspace:7, kind:RawConst> After: %p = addrspacecast ptr addrspace(7) %mapped to ptr insertvalue { ptr, i64 } undef, ptr %p, 0 call { ptr, i64 } @bounce_cluster_pair(...) PTX -O2: mapa.shared::cluster.u64 ; ld.shared::cluster.b32 PTX NO_OPT: mapa.shared::cluster.u64 ; cvta.shared::cluster.u64 ; call.uni bounce_cluster_pair ; ld.b32 ``` ```rust #[inline(never)] #[device] fn bounce_cluster_pair(pair: (*const u32, u64)) -> (*const u32, u64) { pair } let neighbor_ptr = unsafe { cluster::map_shared_rank(addr_of!(SHMEM) as *const u32, neighbor_rank) }; let pair = bounce_cluster_pair((neighbor_ptr, neighbor_rank as u64)); let bounced = unsafe { *pair.0 }; // reads the neighbor block's shared memory ``` ## What we fixed - `AggregateKind::Tuple` was the only aggregate arm that built `MirConstructTupleOp` straight from the translated operands. The verifier requires operand types equal to the element types, so any tuple holding a pointer outside the generic address space (cluster-shared `addrspace(7)`, CTA-shared `addrspace(3)`) failed. - The arm now reads the translated `MirTupleType` element types and runs `cast_to_expected_pointer_type_if_needed` per element, keeping the pointer kind (`RawConst` stays `RawConst`) and declining kind changes as before. - The cluster example round-trips the mapped pointer through a `(ptr, rank)` tuple and a device fn, and its IR check requires the tuple-ABI call and the `addrspacecast` in the pre-opt IR; the compile-only example lane runs that check. ## Address-space contract - A mapped pointer stored in an aggregate (struct, enum, array, and now tuple) is held in the generic form, which is also what CUDA C++ `map_shared_rank` returns. A later deref through the aggregate is a generic access unless LLVM can prove the space; at -O2 it restored `ld`/`st.shared::cluster` in every form tested, and a true mixed cluster/CTA join produced a correct generic load. - A plain deref of the mapped pointer itself still selects `ld`/`st.shared::cluster`. <details><summary>Implementation</summary> - `mir-importer` `translator/rvalue/aggregate.rs`, tuple arm: downcast the translated tuple type to `MirTupleType`, zip `get_types()` with the element values, cast each through the shared helper and thread the previous op. The downcast cannot fail: `translate_type` always yields `MirTupleType` for a Rust tuple. - `translator/rvalue/coerce.rs`: a helper-contract test that an `addrspace(7)` `RawConst` pointer normalizes to the generic `RawConst` type through `mir.cast`. It documents the helper; the regression for the arm is the example plus its IR check. - The cast is `mir.cast <PtrToPtr>`; `mir-lower` `emit_pointer_cast` already emits `llvm.addrspacecast` when the spaces differ, so no lowering change. - cluster example: `bounce_cluster_pair`, the tuple round trip and comparison against the direct read; `verify-code-shape.sh` asserts the call and the `addrspacecast` in `cluster.ll`, read before `opt -O2` because the private helper is inlined there. </details> ## Verification Checked at `95133c9c` merged with `main` (`bc46b1ff`, nightly-2026-10-11): - ✓ Fails before: the PR's cluster source built against main's backend reports the NVIDIA#1453 `MirConstructTupleOp` mismatch; the same source builds at this head. - ✓ `cargo test -p mir-importer`: 1164 passed at the PR head, 1169 on the merged tree, 0 failed. `cargo clippy -p mir-importer --all-targets -- -D warnings` and rustdoc with warnings denied pass; the example's own clippy passes. - ✓ `cargo oxide run cluster` on an RTX 5060 Laptop (author) and an RTX 5090 (sm_120): ring exchange 1001, 1002, 1003, 1000 and mapped store 3003, 3000, 3001, 3002; memcheck and synccheck 0 errors. - ✓ `CUDA_OXIDE_NO_OPT=1 cargo oxide run cluster`: same values with the PTX keeping `cvta.shared::cluster.u64`, `call.uni bounce_cluster_pair` and a generic `ld.b32` through the tuple; memcheck 0 errors. - ✓ `cargo oxide run cluster --emit-nvvm-ir --arch sm_120` (libNVVM): passes. - ✓ `cargo oxide build cluster --arch sm_90` plus `verify-code-shape.sh`: PASS, the compile-only lane's form. - ✓ Further tuple forms (nested, by reference, array of tuples, `Option<*const u32>`, destructured return, `*mut` store through a tuple, layout-reordered `(u8, *const u32, u16)`, control-flow join, pointer equality across the cast, mixed `addrspace(7)`/`addrspace(3)` join) compile and run correctly at -O2 and NO_OPT; memcheck and synccheck 0 errors. A by-value closure capture of the mapped pointer already compiled on `main`. - ✓ PTX and IR of the unchanged tuple examples are identical to `main`. - ✓ Example lockfiles, all repository guards, `git diff --check` and the per-example rustfmt check pass; DCO sign-off present. ## Caveats - At -O2 the example's `bounced == direct` comparison folds to true because LLVM inlines the helper and merges the two loads; the enforced gate is the IR check, and the generic load through the tuple executes only in the NO_OPT build. - Pre-existing and outside this PR: assigning such a pointer into a tuple or struct field after construction (`pair.0 = p`), or `[p, p]` / `[p; 2]` into an addressable local, fails `MirStoreOp` verification on `main` as well; the store path lacks the same normalization. Tracked separately. - No sm_90 hardware run; the PTX ISA defines `cvta.shared::cluster` for sm_90+ and the runs were on sm_120. The legacy CUDA 12 NVVM route at sm_90 fails on this host for `main` too (PTX JIT error) and was not validated. Closes NVIDIA#1453. --------- Signed-off-by: Raul Estrada <raulestradaa@gmail.com> Signed-off-by: nihalpasham <nihalp@nvidia.com> Co-authored-by: nihalpasham <nihalp@nvidia.com>
## What this adds
`#[unroll]` on `for i in a..b` and `for i in (a..b).rev()` now unrolls.
After inlining, such a loop exits through a match on the `Option`
returned by `Range::next`, so the unroll pass found no header exit test
and no induction variable; arrays indexed by `i` stayed in local memory.
A new dialect-mir canonicalization threads that match into a header exit
test, and the signed-counter `checked_add`/`checked_sub` step core emits
is folded to a plain add/sub where the header test proves it cannot
overflow. Casts between identical dialect types (`isize`/`i64`,
`usize`/`u64`) are elided in the importer so `isize` loops look like
`i64` loops.
```text
nested_lane_range (16x16 over a 16-lane accumulator), PTX:
before: 17 ld.local, 17 st.local, .local .b8 __local_depot[64]
after: 0 ld.local, 0 st.local, no depot
```
```rust
let mut acc = [0u32; 16];
#[unroll]
for i in 0..16 {
#[unroll]
for j in 0..16 { acc[j] = acc[j].wrapping_add(mix(input[i * 16 + j])); }
}
```
## What we fixed
- Range `for` loops over any integer counter, including `isize`, with
full `#[unroll]` and partial `#[unroll(N)]`, and `(a..b).rev()` with
full unroll.
- The signed step: core's `checked_add` overflow branch is dead under
the header's strict `lo < hi` test and is folded away; the fold refuses
any other step or an overflow edge that re-enters the loop.
- No-op casts no longer hide the counter from the pass.
- `a..=b`, `.step_by(k)` and custom iterators still warn and stay loops;
the warning text now distinguishes an unrecognized step from a loop that
is not an `a..b` range.
<details><summary>Implementation</summary>
- `canonicalize.rs`: `thread_iterator_exit_test` runs only for a hinted
loop whose header `cond_br` has both successors inside the loop. It
requires single-entry arms ending in `goto` to a two-predecessor join
that reads the discriminant of one argument, each arm building its value
with `mir.construct_enum`, and evaluates the match blocks with a small
evaluator that accepts only storage markers, constants, the discriminant
read, `IntToInt` casts and `eq`/`ne`. Replacements are planned before
any mutation; a match value used where both targets reach is refused.
Dead blocks are removed by the next `simplify_cfg` round in `unroll.rs`.
- `canonicalize.rs`: `fold_checked_step` matches core's
`checked_add`/`checked_sub` by constant 1 on the counter the header
orders strictly, with the overflow edge leaving the loop.
- `analyses/induction.rs`: `match_cmp`, `unwrap_not`, `const_i128` and
`CmpPred::negate` made `pub(crate)` for reuse.
- `mir-importer` `translator/rvalue/expr.rs`: a cast whose operand type
equals the result type returns the operand and emits nothing, beside the
existing reborrow shortcut. Pointer-kind, signedness and width changes
still emit `mir.cast`.
- `unroll_smoke`: new kernels `range_for_loops`,
`signed_range_for_loops`, `signed_partial_range_for`, `isize_range_for`,
`nested_lane_range`; `verify-code-shape.sh` fails the build when
`nested_lane_range` or `isize_range_for` uses local memory.
`nested_lane_range` keeps a runtime inner `while` so LLVM cannot unroll
it alone; it is the discriminating gate. `isize_range_for` is a control,
since LLVM already unrolls it on main.
- Docs: `#[unroll]` macro docs, `thread.rs`, supported-features and the
three book pages list the same supported forms.
</details>
## Verification
Checked at `4d9b8295`:
- ✅ `cargo test -p mir-transforms`: 91 passed. New tests fail on main;
i8/u8 sweeps over every boundary pair, interpreter over every end value
for factors 2, 3 and 4, negative cases for a re-entering overflow edge,
a step of 2 and a value read after both exits.
- ✅ `just test`: 3986 passed; `just book` builds; fmt, clippy, rustdoc
and all 18 `check-*.sh` guards pass; `sync-example-locks.sh --check` OK
after merging current main.
- ✅ `cargo oxide run unroll_smoke` on RTX 5090 (sm_120), LLVM and
libNVVM routes: PASS; memcheck and synccheck 0 errors.
- ✅ `nested_lane_range` with main's backend: 17 `ld.local`, 17
`st.local`, 64-byte depot; here 0/0/none. The 17 pre-existing kernels
are instruction-identical to main.
- ✅ Scratch kernels: u8/i8/u16/i64/u32::MAX boundaries, empty ranges,
negative `isize` bounds, `rev()`, partial remainders for several `n`,
`break` with value, `continue`, runtime bounds, mixed casts; all
correct, 0 local memory, memcheck 0 errors.
## Caveats
- `#[unroll(N)]` on `(a..b).rev()` or on a loop with a runtime start
still warns, as it does for counted `while` loops.
- `a..=b` and `.step_by(k)` are not recognized.
- The fold matches core's current step sequence; a future change there
turns into a warning, which `nested_lane_range` in the compile-only CI
lane would surface.
- `#[unroll(N)]` still emits no loop metadata (NVIDIA#1375), so NVVM may
unroll the remainder again.
- NVIDIA#1494 (const-generic bounds) and NVIDIA#1478 (clone budget) are not
addressed.
---------
Signed-off-by: midagedev <midagedev@gmail.com>
Signed-off-by: nihalpasham <nihalp@nvidia.com>
Co-authored-by: nihalpasham <nihalp@nvidia.com>
## Summary Remove four unused declarations from cuda-oxide's workspace dependency table. The newer Cargo introduced by NVIDIA#1482 reports these entries whenever a workspace command runs. ## Changes Remove `cuda-bindings`, `futures`, `fuzzer`, and `mir-importer` from `[workspace.dependencies]`. No member inherits these declarations. Workspace members and dependencies declared directly by consumers stay intact; neither lockfile changes. ## Testing - Full `cargo metadata --locked --offline --format-version 1` output is identical before and after the change on both `main` and `exp/cuteir-translation`. - `cargo oxide run gemm_views` passed on the experimental checkout with this same four-line cleanup: GPU results matched the CPU reference, and safe/raw pairs were bitwise identical. Its log contains no warnings after also uninstalling the obsolete global `cargo-oxide` executable on that machine. - `git diff --check` passed. Full compiler tests were not rerun for removal of uninherited manifest entries. ## Checklist - [x] Commit signed and signed off. - [x] No new source files or examples; SPDX additions are not applicable. Signed-off-by: nihalpasham <nihalp@nvidia.com>
## Summary Move both cuda-oxide compiler workspaces from the Pliron Git pin to the published 0.19.0 crates. Exact version pins make the dependency reproducible while retaining the configuration without the LLVM FFI bridge. This change targets `main` and applies to the existing compiler and examples. ## Changes - Pin `pliron`, `pliron-derive`, and `pliron-llvm` to `=0.19.0`, update `combine`, both lockfiles, and dependency source/license metadata. - Implement Pliron's new branch operand-range and segment interfaces. Validate malformed metadata, including overflowing counts, before accessing operand ranges or allocating export state. - Adapt printing, source locations, and attribute-map access. Fix function parsing to preserve exactly one body region and round-trip optional attributes. - Add regressions for invalid branches, successor mutations, exporter error handling, and function parsing. ## Testing - [x] `just check` passed from `cuda-oxide/` on the pinned nightly: formatting, strict clippy, tests, guards, intrinsics, and rustdoc/doctests, including the separate backend workspace. - [x] `cargo oxide run vecadd --arch sm_120` passed on RTX 5090: all 1,024 elements correct. - [x] `cargo oxide run array_for_loop --arch sm_120` passed on RTX 5090: scalar and struct array iteration plus generated/mapped array helpers across 32 threads. - Independent review covered the main-based migration, malformed branch metadata, parser invariants, and dependency resolution. - New example: not applicable. ## Checklist - [x] Commit signed and signed off. - [x] Existing SPDX headers preserved; no new source files. --------- Signed-off-by: nihalpasham <nihalp@nvidia.com>
…VIDIA#1332) Type the unclaimed payload bytes of a niche-encoded enum as the widest layout-neutral integer array instead of [N x i8], so the block-argument scalarizer stays within its 16-leaf budget, the Some/None join splits into scalar phis and the pointer behind the filler keeps its global address space (ld.global/st.global instead of generic accesses). Closes NVIDIA#1331. Signed-off-by: Yunus Mutlu <mutluyunus708@gmail.com> Signed-off-by: nihalpasham <nihalp@nvidia.com>
Store shared-memory pointers held in unions as generic pointers and convert at the insert and extract boundaries, so the union's byte image is the same on every target (under the modern NVVM p3:32 layout the raw shared pointer was a 32-bit store read back through an 8-byte field). In-place addresses of such fields fail closed, as the enum payload path already does. Closes NVIDIA#1334. Signed-off-by: Yunus Mutlu <mutluyunus708@gmail.com> Signed-off-by: nihalpasham <nihalp@nvidia.com>
…A#1343) Let a from-end ConstantIndex that follows a slice Subslice lower in both MIR place walkers by taking the runtime length from the rebuilt subslice, so nested slice patterns such as [_, middle @ .., _] followed by [.., last] compile instead of stopping with an unsupported-construct error. Closes NVIDIA#1342. Signed-off-by: Yunus Mutlu <mutluyunus708@gmail.com> Signed-off-by: nihalpasham <nihalp@nvidia.com>
…IDIA#1350) Resolve the input, --output and --decisions-json paths before either write and refuse a collision, including ./ and ../ aliases, hard links (device and inode on Unix), dangling symlinks and a directory as the report path, so the JSON report can no longer silently replace the PTX while the run reports success. In-place PTX rewrite with a separate report stays allowed. Closes NVIDIA#1318. Signed-off-by: YeonwooSung <neos960518@gmail.com> Signed-off-by: nihalpasham <nihalp@nvidia.com>
…IA#1351) TmemGuard::alloc_by and dealloc_by gated the warp-collective tcgen05.alloc and tcgen05.dealloc on threadIdx.x / 32, which names a hardware warp only in a 1D block. Select by the block-linear warp id instead, through one shared helper that the cooperative-groups block reductions also use; 1D blocks are unchanged. Signed-off-by: YeonwooSung <neos960518@gmail.com> Signed-off-by: nihalpasham <nihalp@nvidia.com>
## Summary Update cuda-oxide's compiler workspaces to published Pliron 0.20.0, establishing the dependency baseline for the upcoming CuTeIR interop work. The exporter migration remains on `exp/cuteir-translation`. ## Changes - Pin `pliron`, `pliron-derive`, and `pliron-llvm` to `=0.20.0` across the SIMT and standalone backend workspaces; refresh both lockfiles for the release's required dependency versions. Both graphs resolve one registry copy of each Pliron crate, with `llvm-sys` disabled. - Adapt to the removed `Located` trait, borrowed operation locations, and parser `StreamLocation` trait. Keep owned locations where lowering or debug emission needs them, and preserve the existing location assertions. - Update intrinsic templates and regenerate their outputs; update the dependency inventory and release-specific book reference. ## Testing Validated from current main `51c737fdbd5da9043bdf13bce5b63a777fb47a5c` with the existing `nightly-2026-10-11` pin. - [x] All `just check` gates passed, run individually: formatting, strict workspace/backend clippy, unit and toolkit tests, repository/dependency guards, generated intrinsic checks, and workspace/backend rustdoc. The main unit suite and backend tests used `--locked`; intrinsic ABI history was checked against the base SHA above. - [x] 5,539 tests passed, 0 failed, 11 ignored; 33 doctests passed, 157 ignored. `just book` also passed. - [x] `vecadd` and `array_for_loop` passed on RTX 5090 (`sm_120`). Memcheck with full leak checking reported 0 errors and 0 leaked bytes for both. Reproduce the local gates and GPU checks from `cuda-oxide/`: ```bash just check book cargo oxide run vecadd --arch sm_120 cargo oxide run array_for_loop --arch sm_120 cargo oxide sanitize vecadd --tool memcheck --arch sm_120 -- --leak-check full cargo oxide sanitize array_for_loop --tool memcheck --arch sm_120 -- --leak-check full ``` ## Checklist - [x] Commit signed and signed off. - [x] SPDX checks pass; no new source files. Signed-off-by: nihalpasham <nihalp@nvidia.com>
…1545) ## What this adds A Rust device function can now be called from inside a Tile kernel. `#[tile_export]` marks the function, and `cargo oxide tilelib` packages the crate as a `.tilelib`, the static tile library format tileiras links with `-l`. ```text before: a Tile kernel and a Rust kernel could only share a stream and buffers after: a Tile kernel calls the Rust function in its body, one invocation per tile element ``` ```rust #[tile_export(tile_shape = 128)] pub fn relu(x: f32) -> f32 { if x > 0.0 { x } else { 0.0 } } ``` ```bash cargo oxide tilelib tilelib_relu --arch sm_120 # -> tilelib_relu.tilelib ```  cuda-oxide is the SIMT provider: it ships the code (LTOIR from the existing pipeline) and two small descriptions (the interface and the implementation, as Tile IR bytecode). tileiras does the matching, marshalling and linking. The Tile side only emits the call. Everything lives under `cuda-oxide/`; `cutile-rs` is untouched. ## The contract - `in` scalars are passed by value, `&T`, `&mut T` and fixed arrays as pointers, and a scalar return value becomes a trailing `out` pointer. `#[tile_export]` generates that wrapper next to the function. - `sync = "warp"` means 32 lanes call the function together, so warp shuffles are legal inside it. - Preview content is read only at the assembler's current bytecode version. The command asks the installed tileiras and writes the descriptions at that version: 13.5 for a CUDA 13.5 assembler, the newer version a newer one names, and 13.5 when none can read them yet. <details> <summary>Implementation</summary> - `oxide-tilebc` is a new crate that writes the two descriptions as Tile IR bytecode (`preview$func extern` and `preview$extern_impl`, with the fragment types and layouts they carry). It has a typed model, a structural verifier with the Tile compiler's rules, and a textual printer in the upstream form for the `--dump` output. It writes nothing else of the format; ordinary kernels are not touched. - `oxide-tile-export` is the shared description vocabulary and its rules (one dimension, one footprint per export, written parameters as one contiguous run, whole warps for warp sync). - `#[tile_export]` generates the `extern "C"` wrapper plus a host-only metadata function whose body is the JSON description; the backend reads it from MIR, roots the wrapper, and writes `<crate>.tileexports.json` beside the device artifact. - `cargo oxide tilelib` runs emit-ltoir, checks that the NVVM IR defines every symbol with the parameters the ABI predicts, writes both bytecode halves, and calls `tileiras --preview-create-tile-lib`. On an older assembler it stops after the descriptions with exit 2, or exit 0 with `--library-optional`, which is how the smoketest's new `tilelib` lane runs it. - The `tilelib_relu` example has three consumers: a host program that proves the ABI on CUDA 13.4 (a CUDA C++ caller generated from the descriptions, linked with nvJitLink), a Tile C++ kernel (`cpp/tile_caller.cu`), and a cuTile Python kernel (`python/call_from_cutile.py`). </details> ## Verification Checked at `a7c465e9d`: - ✅ oxide-tilebc 15 tests (hand-built byte tables, verifier rules, version parsing), clippy with warnings as errors, rustdoc. - ✅ cargo-oxide 255 tests, cuda-macros 132 tests plus the trybuild pass and nine compile-fail cases, oxide-tile-export 16 tests; all 18 repository guards, lockfile sync, cargo deny, book with warnings as errors. - ✅ RTX 5090: `cargo oxide run tilelib_relu` verifies all three exports through the native ABI on CUDA 13.4. - ✅ RTX 5090 with a 13.5-capable tileiras: `cargo oxide tilelib` builds `tilelib_relu.tilelib`, and the Tile C++ kernel calls all three exports through it (3 of 3, bit for bit). The decoded descriptions equal the accepted reference bytecode. - ✅ cuTile Python 1.7 dry run: the kernel compiles to 13.5 bytecode carrying the `preview$call` to `relu`. ## Caveats - The `.tilelib` itself needs a CUDA 13.5 or newer tileiras; on 13.4 the command writes the descriptions and the LTOIR and says so. - The cuTile Python run needs a cuTile Python that lists the assembler's bytecode version; the Tile C++ path has no such gap. - One-dimensional tiles, `f16` but not `bf16`, concrete shapes only. Linkage providers (implementations compiled on demand) are not part of this change. --------- Signed-off-by: nihalpasham <nihalp@nvidia.com>
…DIA#1493) Add a "Compiler ordering around the barrier" section to the shared-memory chapter: how each backend emits sync_threads, a table saying per access form whether the compiler keeps it on its side of the barrier, and short notes on raw pointers, cp.async and references. The ptx_asm! row of the feature list links to it. Signed-off-by: Alejandro Soto Franco <sotofranco.eng@gmail.com> Signed-off-by: nihalpasham <nihalp@nvidia.com>
…VIDIA#1329) The installer refreshed the shared backend at ~/.cargo/cuda-oxide/librustc_codegen_cuda.so with fs::copy, which truncates and rewrites the inode that every running rustc has mapped, so a reinstall during a kernel build killed that rustc. Copy the new build to a sibling temporary and rename it over the name; clear the commit record just before the rename so a failed copy keeps the old backend and its record; let `cargo oxide update` remove a temporary left by a killed install. Signed-off-by: midagedev <midagedev@gmail.com> Signed-off-by: nihalpasham <nihalp@nvidia.com>
The drop no-op proof erased a destructor call it never proved returns, so a destructor that loops forever or recurses into itself was deleted and the kernel ran on past it. The proof is now a three-state depth-first walk: a reachable cycle, or a call back into an instance already under proof, keeps the call; a reconverging if/else destructor is still erased. The drop_glue example asserts the decision in the pre-opt IR. Signed-off-by: Yunus Mutlu <mutluyunus708@gmail.com> Signed-off-by: nihalpasham <nihalp@nvidia.com>
A one-page figure and an "Inlining helpers" subsection on the device-functions page for the inline intent mapping that NVIDIA#1511 landed: #[inline] as inlinehint, #[inline(always)] as alwaysinline, #[inline(never)] as noinline, and the collector-boundary exception for generic #[device] helpers. Every quoted line comes from the inline_intent example built on an RTX 5090. The inline_intent lockfile committed with NVIDIA#1511 predates the NVIDIA#1545 re-alignment, so the first build of the example added four entries and left the tree dirty; a resolve against the existing lock adds them at the versions the other 239 locks carry. Signed-off-by: nihalpasham <nihalp@nvidia.com>
…IDIA#1336) Add cp_async_bulk_tensor_{1..5}d_g2s_cta, TMA tile loads through the PTX .shared::cta destination form, and SharedPtr<T>, a pointer into the issuing CTA's shared memory whose address space is part of its type. SharedPtr lowers to ptr addrspace(3) across helper parameters, returns, struct fields, generics and closures, so the CTA form never converts a generic address, and ptxas on sm_120a keeps setmaxnreg instead of inserting the runtime CTA check that dropped it. The cluster family is unchanged. The five catalog records carry ABI ids i1034 to i1038 behind overlay shard schema 65 with evidence recorded under the pinned LLVM and libNVVM profiles; the exporter refuses shared or local pointers inside aggregate kernel parameters; SharedArray::shared_ptr is unsafe like as_raw_mut_ptr; shared_memory_helper, shared_ptr_api and error_kernel_shared_ptr_param cover the type across call boundaries; the book documents the type, its limits and the CTA form. Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com> Signed-off-by: nihalpasham <nihalp@nvidia.com>
Move the FP16 m16n8k8 MMA floors to PTX 6.5 and sm_75 at their one source, the recipe table, so the policy, both backend lowerings, the catalog, the generated target requirements and the stub docs follow by regeneration. BF16 and TF32 m16n8k8 and FP16 m16n8k16 keep PTX 7.0 and sm_80, as the PTX ISA says and ptxas enforces. The Turing evidence is recorded under the current LLVM profile from real llc, ptxas and opt runs, the generated summaries no longer call the Turing forms Ampere forms, and the codegen negative test covers every sm_80 member of the admission at sm_75. Signed-off-by: Praveen Perera <me@praveenperera.com> Signed-off-by: nihalpasham <nihalp@nvidia.com>
Adds an NVFP4 MoE benchmark that validates outputs and measures CUDA graph replay latency for power-of-two batch sizes from 1 through 256 using synthetic inputs. Performance optimization is planned for a future version. Validated on Blackwell with CUDA 13.3: full batch sweep, formatting, clippy, CPU/GPU suites, 27 supported examples, and eight benchmark smoke tests (smoke-test flag corrected locally in the runner). Signed-off-by: Melih Elibol <elibol@users.noreply.github.com>
… the toolkit ships (NVIDIA#1555) ## What this adds The tile-library chapter leaves the published book until the CUDA toolkit that carries the feature is announced. The code stays. ```text before: the book describes #[tile_export] and cargo oxide tilelib, which need a toolkit that is not out yet after: the chapter, its figure and every mention are out of the book and the READMEs; the command is hidden from --help ``` ## What we fixed - Removed `advanced/tile-libraries.md`, its figure and the table-of-contents entry. - The supported-features row for Tile interop keeps the inter-kernel story and points at the tracking issue only. - The crate tables in the README and the architecture page keep their two rows, since the inventory guard requires every member, but describe the crates without the feature. - `cargo oxide tilelib` is hidden from `--help`, which is what lets the command-reference guard accept the missing row. The command still runs. ## Verification Checked at `1b2011d44`: - ✅ cargo-oxide 268 tests, formatting, book with warnings as errors. - ✅ Guards: CLI documentation coverage (16 visible subcommands, 3 hidden), crate inventory, source references, book API names, book catalog stamp, book import order, host API paths, SPDX, test matrix, smoketest contract, reserved prefixes, error example status. ## Caveats - The example `tilelib_relu` and its README stay in the tree; they are not part of the book. - Reverting this PR restores the chapter as it was. Signed-off-by: nihalpasham <nihalp@nvidia.com>
…sembler reads it (NVIDIA#1556) ## What this adds A `.tilelib` written by `cargo oxide tilelib` now loads on the CUDA 13.5 assembler. ```text before: 13.5 tileiras rejects the implementation descriptor: "concrete shape cache rank (1) does not match symbolic shape rank (0)" after: the descriptor parses, and a Tile C++ kernel calls the three exports through the library ``` ## What we fixed - The 13.5 assembler keeps a full-rank symbolic shape on `preview$fragment`, one integer expression per dimension. The compiler trunk switched on 2026-09-08 to holding only the dynamic dimensions, and the writer produced that trunk form for every version. - The fragment type now encodes per version: full rank below 13.6, dynamic-only from 13.6 on. Nothing else in the descriptors changes. <details> <summary>Implementation</summary> - `oxide-tilebc`'s type table carries the bytecode version, and the fragment arm branches on it. - The 13.5 golden in the writer tests now holds a one-element symbolic tuple. </details> ## Verification Checked at `514be3a18`: - ✅ oxide-tilebc 15 unit tests, clippy with warnings as errors, rustdoc. - ✅ CUDA 13.5.40 release toolkit, RTX 5090: `cargo oxide tilelib tilelib_relu --arch sm_120` builds the library, its disassembler reads both descriptors, and `cpp/tile_caller.cu` passes 3 of 3 through it. - ✅ The 13.6 development build still reads the dynamic-only form: the same library built at 13.6 passes 3 of 3 there too. ## Caveats - The version switch keys on 13.6 because that is where the trunk form first ships; if a 13.5 update moves to the trunk form, the boundary moves with it. Signed-off-by: nihalpasham <nihalp@nvidia.com>
## Summary Pinned the third-party actions to commit SHAs. Closes NVIDIA#1550 ## Checklist - [x] All commits signed off (`git commit -s`) Signed-off-by: Jacob Wujciak-Jens <jwujciakjens@nvidia.com>
…NVIDIA#1602) Preserve source locations and inline caller chains for const-generic arrays, generated reduction bodies, and checks relocated to the host. Keep deduplicated check origins private to the compiler and render them through the existing diagnostic API. Part of NVIDIA#1475. The remaining frontend items and debugger/backend fixes remain tracked separately; this does not close the full source-location audit. Validated with the CPU/assembler suites, source-location regressions, live CUDA-GDB tests, GPU host-check enforcement, host/Tile semver checks, and all 25 active CI checks on 825de3e. Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Advance the shared host and Tile crates and their path dependency constraints to v0.5.0. Update local package versions in the affected lockfiles without changing registry dependencies. Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
…NVIDIA#1598) ## What this adds A preview of CuTe kernels in Rust, under `cuda-oxide/cute-rs-preview`. Layouts, swizzles, operand roles, TMA pipelines and tensor-core MMA are Rust types; the compiler decodes them, verifies the whole module, and hands CUTLASS MLIR to the pinned CUTLASS 4.7 compiler. Two targets, `sm_100a` and `sm_120a`; everything else is rejected before compilation. ```text before CuTe kernels need Python: @cute.jit traces live MLIR wrappers at run time after cargo oxide run blockscale_gemm_cute --arch sm_120a Rust types -> dialect-cute -> whole-module verification -> CUTLASS 4.7 -> cubin in the binary ``` ```rust type ASmem = Composed<Swizzle<3, 3, 3>, 0, RowMajor<128, 64>>; #[kernel] pub fn gemm(#[grid_constant] a: &TmaDesc<f16, ASmem>, m: u64, n: u64, k: u64) { // the same ASmem type encodes the host descriptor and checks the device copy } ```  - Four crates: `cute-rs` (tensors, tiles, copies, pipelines, tiled MMA), `ctm-rs` (typed instruction primitives: TMA, mbarrier, tcgen05, TMEM, mma.sync), `cute-layout` (layout algebra and compile-time proofs) and `dialect-cute` (the Pliron dialect and verifier). - Compiler side: `mir-importer` recognises `cute-rs` calls by definition path and emits `dialect-cute`; `cuda-oxide-mlir-export` maps MIR, NVVM and CuTe operations to CUTLASS MLIR through the published `pliron-mlir-interop` printer; `cuda-oxide-codegen` drives the pinned `libCutlassCompiler.so` and embeds the validated cubin. `cargo oxide toolchain install cutlass` installs the compiler and checks its digests. - 14 example crates (10 on `sm_120a`, 4 on `sm_100a`) and two end-to-end models on `sm_120a`: Qwen3-TTS 1.7B and Qwen3-8B in BF16 and NVFP4. - Four new intrinsics: the `m16n8k64` MXF4 register MMA and three relaxed cluster barrier operations. - Book chapter "CuTe in Rust (preview)" with a support matrix against CuTeDSL 4.7, build and contracts pages. The preview covers a validated subset of CuTeDSL 4.7 on the two targets, not feature parity: one of eleven tcgen05 MMA kinds, three of seven warp-level kinds, 32x32b TMEM only, fixed epilogues, a static scheduler. The support matrix page lists every row with what the preview has. ## What we fixed - The CuTe stack folder is named `cute-rs-preview` and the word preview appears where readers look: the folder README, the book chapter and index card, the four crate descriptions. - Two book snippets quoted a producer method that does not exist; they now show the acquire, elected expect and barrier sequence the kernel uses. - The elementwise README selected no backend and no target; the block-scaled GEMM README claimed `sm_121`. - The stack's docs folder (compiler history, a superseded Pliron 0.19 experiment, duplicated figures) is gone; two pages moved into the book and one paragraph into the mlir-export README. <details><summary>Implementation</summary> - `cuda-oxide-mlir-export/src/profile.rs` holds the target gate: `CutlassFullCuteMlir22::new` accepts `sm_100a` or `sm_120a`; `tcgen05` operations are refused on `sm_120a` and MXF4 register MMA off `sm_120a`. The primitives-only profile `cutlass-primitives-mlir22` has the same gate. - `dialect-cute/src/verify/` runs before any MLIR is built: tensor provenance, pipeline lifecycle, cursor provenance, scheduler grids, MMA placement, epilogue phases, and a ghost-escape audit so compiler-only types never reach runtime storage. - `mir-importer/src/translator/cute/` is a closed world: 88 recognised entry points over 60 operation kinds, 15 audited source-abstraction modules; any other `cute_rs::` call is an error. - The merge with `main` renumbered the four new ABI ledger entries to `i1039` to `i1042` behind main's `i1030` to `i1038`, kept main's lockfiles and let cargo add the new crates, and regenerated every catalog-derived file; `check`, `probe --all` and `check-abi-history --base-ref main` pass. - The folder rename is one scripted commit; package names did not change, so `cute_rs::` definition paths are unchanged. </details> ## Verification Checked at `56384d465`: - ✅ `just fmt-check`, `just clippy`, `just test`, `just test-cuda`, `just check-guards`, `just doc-check` and `just check-intrinsics origin/main` pass on the merged tree. - ✅ `cargo test --workspace --all-targets --no-fail-fast` over 115 test binaries, `cargo test -p oxide-artifacts --features object`, and the nested `rustc-codegen-cuda` workspace pass. One pre-existing race, `ptx::tests::debug_builds_fail_when_opt_drops_the_debug_graph`, failed once under parallel load (a just-written temp file reported busy) and passed three times alone. - ✅ The generator gates: outputs current, 2,427 probe configurations across 1,042 routes, ledger append-only against `main` with 1,042 entries. - ✅ `scripts/sync-example-locks.sh --check` after aligning the 13 cute example lockfiles with the standard examples. - ✅ `scripts/check-primitives-api-compile.py` for the `api`, `fp16` and `blockscale` fixtures, and the 13 CUTLASS consumer tests with the pinned 4.7.0 library. - ✅ The book builds with warnings as errors; 436 relative links resolve; the source-reference, host-API-path, SPDX and crate-inventory guards pass. - ✅ On an RTX 5090 through CUTLASS 4.7: `elementwise_cute`, `rmsnorm_cute`, `swiglu_cute`, `ffn_cute`, `rope_cache_cute`, `blockscale_gemm_cute`, `blockscale_gemm_ctm`, `nvfp4_gemv_cute`, `fmha_prefill_sm120a_cute` and `fmha_decode_sm120a_cute` run and pass their reference comparisons; the FMHA pair also times against CuTeDSL 4.7.0 (prefill 18.0 us against 23.1 us, decode 4.4 us against 12.0 us). The four `sm_100a` examples compile. `vecadd`, `mma_mxf8f6f4` and `cluster` run on the native path. - ⏳ Not run locally: `scripts/smoketest.sh --compile-only`, which publishes to the shared backend cache; the examples-compile lane covers it here. ## Caveats - Two targets only. `sm_100f`, `sm_103a`, `sm_110a`, `sm_120f`, `sm_121a`, Hopper and Ampere are rejected at the profile gate. - CuTe kernels build only with `CUDA_OXIDE_DEVICE_BACKEND=cutlass-mlir`; the native backend has no CuTe lowering and currently stops with a verification error rather than a dedicated message. - NVFP4 in the Qwen3-8B model runs on CUDA cores through scaled loads and `cute.dot`; the NVFP4 tensor-core atom exists in the dialect and exporter but has no producer yet. - No CI job compiles the cute examples; the unit-test, clippy and fmt lanes cover the crates and the example manifests, and GPU runs are manual. - APIs can change between releases. The compatibility aliases in `cute-rs` and `dialect-cute` stay for now. --------- Signed-off-by: nihalpasham <nihalp@nvidia.com>
NVIDIA#1615) ## Summary Update `cuda-bindings`, `cuda-core`, and `cuda-core-derive` from `0.4.0` to `0.5.0` in the 17 `Cargo.lock` files introduced by NVIDIA#1598 (`primitives_api_compile` plus the 14 `cute-rs-preview` examples and 2 models) to align with the `v0.5.0` bump in NVIDIA#1614 (`0143af5cc`). ```text Before: 17 cute-rs-preview / fixture Cargo.lock files -> cuda-core 0.4.0 (fails --locked on main) After: 17 cute-rs-preview / fixture Cargo.lock files -> cuda-core 0.5.0 ``` ## Validation - `bash cuda-oxide/scripts/sync-example-locks.sh --check` - `bash cuda-oxide/scripts/check-dependency-licenses.sh` - `bash cuda-oxide/scripts/check-example-license-policy.sh` ## Checklist - [x] All commits signed off (`git commit -s`) Signed-off-by: nihalpasham <nihalp@nvidia.com>
## Summary Move `cuda-oxide` from a pinned daily `nightly` toolchain (`nightly-2026-10-11`) to official **stable Rust `1.99.1`**, unifying the entire monorepo (`cuda-core`, `cutile-rs`, and `cuda-oxide`) on stable Rust releases: ```text Before: cuda-oxide/rust-toolchain.toml -> channel = "nightly-2026-10-11" (daily snapshot) After: cuda-oxide/rust-toolchain.toml -> channel = "1.99.1" (official stable release + RUSTC_BOOTSTRAP=1) ```  --- ### Before vs. After - **Before (`nightly` pin):** `cuda-oxide` pinned a moving daily `nightly` snapshot while the root host workspace and `cutile-rs` used stable Rust. Any mid-cycle nightly bump (such as pulling in an upstream `rustc_public` addition) forced an unscheduled compiler upgrade with unrelated daily churn across 1,500+ files. - **After (Stable `1.99.1` pin):** `cuda-oxide` pins official stable Rust `1.99.1`, which has already completed Rust's 12-week nightly and beta stabilization cycle and ships **LLVM 23 (`23.1.1`)** with an NVPTX backend and TableGen intrinsic definitions identical to `nightly-2026-10-11`. ### Why This Works on Official Stable Rust 1. **Stable `rustc` and Nightly `rustc` are the same binary:** Official stable `rustc` builds do not strip out `-Z` flags (`-Z codegen-backend=...`, `-Z mir-enable-passes=-JumpThreading`, `-Z always-encode-mir`) or `#![feature(rustc_private)]`. They gate them behind `RUSTC_BOOTSTRAP=1` (the environment variable `rustc` uses when bootstrapping `core` and `std` on stable). 2. **`rustup` ships `rustc-dev` and `rustc_public` on stable releases:** Running `rustup component add rustc-dev` on stable `1.99.1` provides `librustc_driver.so`, `librustc_public.rlib`, and `librustc_public_bridge.rlib` out of the box. 3. **Production precedent:** This follows the same model used by Microsoft's `rustc_codegen_utc` and Ferrocene (building a `CodegenBackend` against a vetted stable Rust release tag rather than tracking daily nightlies). ### Upstreaming to `rustc_public` on the Stable Cycle  Because `rustc_public` only wraps data that already exists inside `rustc_middle` and `rustc_abi`, `crates/mir-importer/src/translator/facts.rs` acts as our quarantine boundary when we upstream new `rustc_public` queries (such as `rust-lang/rust#163242`): - **Day 1:** Land the feature immediately on the current stable pin by reading the underlying `rustc_middle` / `rustc_target` field in `facts.rs` via `rustc_internal::internal(tcx, ...)`, while upstreaming the `rustc_public` getter in parallel. - **Next scheduled stable release:** Swap the temporary internal read in `facts.rs` to the upstreamed `rustc_public` method call. --- ## Changes in This PR - **Pin `1.99.1` across all SIMT toolchain surfaces:** `cuda-oxide/rust-toolchain.toml`, `cargo oxide new` (`scaffold.rs`), `.devcontainer/devcontainer.json`, `intrinsics/upstream.lock`, and the book/README install commands. - **Configure `RUSTC_BOOTSTRAP=1` automatically:** Set in `cuda-oxide/.cargo/config.toml`, `cuda-oxide/Justfile`, `cargo-oxide` (`backend.rs`, `codegen_env.rs`, `fmt.rs`, `host_cargo.rs`), and the example CI loops so `-Z` compiler flags, `#![feature(rustc_private)]`, `#![feature(f16)]`, and `cargo-features = ["profile-rustflags"]` work transparently on stable `1.99.1`. - **Adapt to `1.99.1` compiler APIs (~12 lines):** - `cuda-host`: Use `core::intrinsics::type_id::<T>()` in `type_id_u128`. - `mir-importer`: Query `tcx.fn_abi_of_instance` inside `facts.rs` (`reference_param_validity`) for typed `ArgAttributes` (`NoAlias`, `NoFree`, `pointee_size`), and match `ValueAbi::Scalar` in `types.rs`. - `rustc-codegen-cuda`: Match `1.99.1`'s `CodegenBackend` (`init`, `target_config`, `CompiledModule`), `ShimKind::FnPtrAddr`, and `layout.peel_transparent_wrappers`. - **Update `cuda-oxide-book`:** Add the toolchain stabilization, `facts.rs` upstreaming workflow, and standalone bundle sections with four SVG diagrams under `cuda-oxide-book/compiler/images/`. ## Validation - **`just check` & `just book`:** `fmt-check`, `clippy` (workspace + backend + all 258 example/model manifests), `test`, `test-cuda`, `check-guards`, `check-intrinsics` (all 1,042 routes / 2,427 target probes on `LLVM version 23.1.1-rust-1.99.1-stable`), `doc-check`, and Sphinx HTML book build (`-W`) all pass. - **SIMT examples (`scripts/smoketest.sh` on RTX 5090):** 240 / 240 pass. - **`cute-rs-preview` examples and models:** 16 / 16 pass (all 12 `sm_120a` workloads execute and pass on RTX 5090; all 4 `sm_100a` workloads compile to validated direct-ABI cubins via CUTLASS 4.7). ## Checklist - [x] All commits signed off (`git commit -s`) - [x] SPDX headers on new source files Signed-off-by: nihalpasham <nihalp@nvidia.com>
… clean up docs (NVIDIA#1616) ## Summary Follow-up cleanup after moving `cuda-oxide` to stable Rust `1.99.1` (NVIDIA#1609) and landing `cute-rs-preview` (NVIDIA#1598): ```text 1. cuda-oxide/flake.lock + flake.nix -> bump rust-overlay for Stable 1.99.1 + set RUSTC_BOOTSTRAP=1 2. cargo oxide new (scaffold.rs) -> write .cargo/config.toml with RUSTC_BOOTSTRAP = "1" 3. check-toolchain-parity.sh -> verify flake.lock's rust-overlay rev has the pinned manifest 4. .github/workflows/unit-tests.yml -> drop unused apt llvm-21 install (sysroot llvm-tools provides LLVM 23) 5. CONTRIBUTING / AGENTS / README / book -> update remaining nightly wording + clean up book backticks & SVG ``` ### Changes - **Nix flake (`cuda-oxide/flake.lock`, `cuda-oxide/flake.nix`):** - Bump `oxalica/rust-overlay` from `c1cfa7a2c8` (`2026-10-11`, two days before `1.99.1` was published) to `5f267b7f32` (`2026-10-22`) so `nix develop` finds `Stable 1.99.1`. - Set `RUSTC_BOOTSTRAP = "1"` in `flake.nix` (`cargoOxideCommonArgs` and `devShells.default.shellHook`). - Use a deterministic per-user directory (`${XDG_RUNTIME_DIR:-/tmp}/cuda-oxide-nvidia-driver-$(id -u)`) for non-NixOS driver symlinks so `nix develop --command` does not leak `/tmp/nix-nvidia-driver.XXXXXX` directories. - **Project scaffold (`crates/cargo-oxide/src/commands/scaffold.rs`):** - Write `.cargo/config.toml` (`[env]\nRUSTC_BOOTSTRAP = "1"\n`) in `cargo oxide new` so plain `cargo check`, `cargo build`, and `rust-analyzer` work out of the box in standalone projects on stable `1.99.1`. - **Toolchain parity guard (`scripts/check-toolchain-parity.sh`):** - Verify via an HTTP `HEAD` check that the `oxalica/rust-overlay` revision locked in `cuda-oxide/flake.lock` actually contains `manifests/stable/<channel>.nix` (or `manifests/nightly/...`), and document the `nix flake update rust-overlay` step in `check-toolchain-parity.sh` and `compiler/rustc-public.md`. - **CI workflow (`.github/workflows/unit-tests.yml`, `.github/workflows/status-guard.yml`):** - Remove the redundant `Install LLVM 21` step on `cuda-oxide-codegen` (`LlvmToolchain::resolve` already selects the Rust sysroot's LLVM 23 `llc`/`opt` from `llvm-tools`). - Update the `status-guard.yml` comments to match the current single SIMT stable pin and `flake.lock` check. - **Contributor, agent, and book docs (`CONTRIBUTING.md`, `AGENTS.md`, `cuda-oxide/README.md`, `cuda-oxide-book/`):** - Replace remaining "nightly toolchain" wording in `CONTRIBUTING.md`, `AGENTS.md` (`CLAUDE.md`), `cuda-oxide/README.md`, `getting-started/installation.md`, and `getting-started/hello-gpu.md`. - Remove backticks from `# ` sidebar page titles, `## ` / `### ` section headings, and prose crate/tool names across the `cute/` chapter and `compiler/rustc-public.md`. - Scale bar widths in `models/qwen3_8b_cute/svg/03-rtx5090-benchmark-comparison.svg` so the right-hand throughput labels fit inside the panel with >=45px margin. ## Validation - `cargo test -p cargo-oxide scaffold` - `bash cuda-oxide/scripts/check-toolchain-parity.sh` - `just fmt-check` - `just check-guards` - `just book` (Sphinx HTML build with `-W`) ## Checklist - [x] All commits signed off (`git commit -s`) Signed-off-by: nihalpasham <nihalp@nvidia.com>
…k for 1.99.1 (NVIDIA#1617) ## Summary Update remaining references to a pinned `nightly` toolchain across the root `README.md`, umbrella `docs/index.md`, `cuda-oxide-book`, crate READMEs, and comments, and fix `cargo oxide doctor` step 1 so it does not reject stable `rustc 1.99.1`: ```text Before: README.md / docs/index.md / crate READMEs say "Pinned nightly"; `cargo oxide doctor` fails step 1 with "expected nightly, got: rustc 1.99.1" After: Docs and comments reflect the pinned stable Rust toolchain (`1.99.1`); `cargo oxide doctor` passes step 1 and verifies the pin in step 2 ``` ### Changes - **Root & umbrella documentation (`README.md`, `docs/index.md`, `Cargo.toml`, `rust-toolchain.toml`):** - Update the "Picking a model" table, SIMT overview, and Contributing section from `Pinned nightly` to `Pinned stable Rust (1.99.1), managed by cargo oxide`. - **`cuda-oxide` book & crate READMEs:** - Update `appendix/building-from-source.md`, `gpu-programming/error-handling-and-debugging.md`, `crates/cargo-oxide/README.md`, `crates/cuda-oxide-codegen/README.md`, and `cute-rs-preview/{cute-rs,ctm-rs,cute-layout}/README.md`. - **`cargo oxide doctor` (`crates/cargo-oxide/src/commands/doctor.rs`, `src/backend.rs`):** - Drop the `version.contains("nightly")` gate in `doctor` step 1 (step 2's `doctor_report_toolchain_pin` already checks that the active toolchain matches `rust-toolchain.toml` and has all required components), and update the mismatch string in `backend.rs`. - **Comments (`cuda-oxide/Cargo.toml`, `rustc-codegen-cuda/Cargo.toml`, `cuda-host/src/lib.rs`, `cuda-oxide-codegen/src/lib.rs`, CI workflows, and scripts):** - Replace leftover "pinned nightly" comments with "pinned toolchain". ## Validation - `cargo run -p cargo-oxide -- doctor` - `cargo test -p cargo-oxide` - `just fmt-check` - `just check-guards` - `just book` ## Checklist - [x] All commits signed off (`git commit -s`) ### Python reference environment (cute-rs examples and models) One setup for every CuTeDSL comparison, the same three lines in each README: ```bash cargo oxide toolchain install cutlass python3 -m venv cute-rs-preview/.venv cute-rs-preview/.venv/bin/pip install -r cute-rs-preview/requirements.txt ``` - **`cute-rs-preview/requirements.txt`:** shared pins for the DSL 4.7.0 family; the per-example `requirements.txt` files include it. The blockscale references keep their own DSL pins. - **Runners:** one lookup order everywhere: `--python` where it exists, `CUDA_OXIDE_CUTE_PYTHON`, the example's `.venv`, the shared `cute-rs-preview/.venv`, then `python3`. The SM100 runners no longer look for the fp16_gemm venv; fp16_gemm's automatic setup now creates the shared environment. - **Fail fast:** each `compare.py` checks for `numpy`, `torch`, `cutlass` and `cuda` before any GPU work and prints the setup commands instead of a traceback after the Rust build. - **READMEs:** every example and model README names its setup inline; the Rust-only ones name `cargo oxide toolchain install cutlass`. The devcontainer installs `python3-venv`. Validation: `cargo fmt --check`, `cargo clippy --all-targets -- -D warnings` and `cargo test` for the three edited example crates; `python3 -m unittest public_reference.test_setup`; the status-guard scripts. --------- Signed-off-by: nihalpasham <nihalp@nvidia.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com>
## Summary Fixes NVIDIA#1618. [RUSTSEC-2026-0386](https://rustsec.org/advisories/RUSTSEC-2026-0386) rejects `tokio` before 1.53.3 on the 1.53 line. These locks move to 1.54.1: - `cuda-oxide/crates/rustc-codegen-cuda/examples/async_mlp` - `cuda-oxide/crates/rustc-codegen-cuda/examples/async_vecadd` - `cuda-oxide/crates/rustc-codegen-cuda/examples/cutile_inter_kernel` - `cutile-rs/Cargo.lock` (`cutile-rs/deny.toml` checks advisories and does not ignore this one) - `cutile-rs/cutile-benchmarks/paper/app_exp1_work_distribution/bimodal_gemms_rust` `launch_overhead_rust` still pins `tokio` 1.53.0. Its manifest patches dependencies to a local path outside this tree, so that lock was not updated here. ## Test plan - [ ] `cargo-deny / deny.toml governs the example workspaces` passes - [ ] `bash cuda-oxide/scripts/sync-example-locks.sh --check` still reports one version per crate - [ ] `cargo deny --manifest-path cutile-rs/Cargo.toml --locked check` passes advisories --------- Signed-off-by: Roman Ivanov <rivanov@nvidia.com>
… diagrams (NVIDIA#1640) ## Summary - Lighten and desaturate the `@media (prefers-color-scheme: dark)` filler colors across the dual-theme SVGs in `cuda-oxide-book/` so amber, purple, teal, coral, and neutral containers render as subtle muted slate tones instead of heavy dark brown/indigo fills. - Lift the dark-mode canvas and inner `.k-box` pill backgrounds so nested pills remain clearly readable against container fills. - Replace two hardcoded `#5c5c5c` text fills in `zipped-divide-and-slice.svg` and `host-device-tma-type-agreement.svg` with the theme-aware `.ts` class. #### Test plan - [x] `just -f cuda-oxide/Justfile book` builds with zero warnings - [x] Verified rendered dark-mode diagrams in the local book preview Signed-off-by: nihalpasham <nihalp@nvidia.com>
A size relation could name only kernel parameters and integer literals. A kernel whose buffer sizes follow a named constant had to write the value in the contract and tie it to the constant with a const assert. An operand may now be a constant: one identifier with no lower-case letter that names no parameter (`TILE`, a const generic `N`), or a path of two or more segments (`P::WIDTH`, `crate::m::TILE`). One lower-case identifier that names no parameter is still the "unknown identifier" error. Validation and code generation use one classifier. The macro cannot see a constant's type, so the check widens it through a sealed trait (u8/u16/u32/u64/usize) instead of an `as` cast; a signed or wider constant is an E0277 at the relation. A kernel outside `#[cuda_module]` names its constants in a never-called function beside it, so a wrong constant fails there too. A contract over parameters and literals expands as before. Tests: pass and compile-fail cases for constants, paths, const generics, and signed or misspelt constants, in and outside a module. Signed-off-by: midagedev <midagedev@gmail.com>
…nels A `#[kernel]` that a `macro_rules!` call declares inside a `#[cuda_module]` expands after the module macro runs, so the module does not see it. Like an `include!`d kernel it gets no launcher, and a call fails with E0599. When every kernel of a module comes from such a call, the error was only "cuda_module found no #[kernel] functions in this module". It now names the call and the two ways out: declare the kernels in the module, or let the macro expand to the whole module. Tests pin the E0599 boundary next to the `include!` one, and the whole-module macro as supported. Signed-off-by: midagedev <midagedev@gmail.com>
When an enum payload is a struct or tuple whose pointer leaf is the niche
carrier (Option<(f32, f32, &mut T)>, Option<{ &mut T, u64, u64 }>), only the
pointer leaf got a slot. The other leaves fell into filler bytes, so
construct wrote them through a spill slot and the payload read loaded them
back at another width. SROA cannot rebuild mixed-type slices, and the slot
reached PTX as a local depot.
Now every leaf of such a payload claims a typed slot at its rustc offset,
reusing the carrier's slot where the types match, and construct/payload
build and read the value with insertvalue/extractvalue. The enum's size,
alignment, field offsets and niche encoding are unchanged; any payload the
rule cannot prove falls back to the old lowering.
Adds the enum_niche_payload example.
Signed-off-by: midagedev <midagedev@gmail.com>
* feat(cargo-oxide): select extra-rustflags by target triple `extra-rustflags` in `.cargo/cuda-oxide.toml` was one flat list, and the `CARGO_ENCODED_RUSTFLAGS` cargo-oxide exports hides Cargo's own `target.<triple>.rustflags`. A project whose host code wants a CPU flag therefore gave it to every host: `-C target-cpu=znver3` written for an x86_64 workstation reached an aarch64 build (DGX Spark), where rustc warns that znver3 is not a processor on every crate and crashes inside LLVM on some. A `[target.<triple>]` table now holds an `extra-rustflags` list for one triple. A build gets the top-level list, then the table of the triple it builds for: the command's `--target` (before any `--`), else `CARGO_BUILD_TARGET`, else the host triple of `rustc -vV` run where Cargo runs. A triple with no table gets the top-level list alone. With tables present and no triple known, or two different `--target` values in one command, the build stops with a named error instead of guessing. A config without tables resolves exactly as before and runs no extra `rustc -vV`. `cfg(...)` keys and fields other than `extra-rustflags` are refused by name, so a table cargo-oxide cannot honor never builds without its flags; `doctor` reports them through the existing config check. Tests: a per-target config built with `--target aarch64-unknown-linux-gnu` gets `-C target-cpu=native` and no znver3, and with `--target x86_64-unknown-linux-gnu` the x86 list (red before this change: the table was ignored and aarch64 got no flags); a flat list still reaches every triple in the same position; selection order, unlisted and unknown triples, the host fallback, `--target` parsing and the refused shapes. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Signed-off-by: Alis Volat Propriis <alisvolatprop12@proton.me> * fix(cargo-oxide): resolve host-tuple and probe the build's own rustc Review of the per-target `extra-rustflags` selection found three gaps: - `--target host-tuple` (or `CARGO_BUILD_TARGET=host-tuple`), which Cargo accepts as the host triple, was looked up as a table named `host-tuple`, so the host's table was silently skipped. It now resolves through the host probe, asked at most once, and counts as the same target as the host's own name. - The host probe ran `rustc` from PATH with the process environment. It now runs the compiler Cargo would run for the command (`RUSTC`, else `CARGO_BUILD_RUSTC`, the command's own value before the inherited one, else `rustc`) with the command's directory and environment changes, so a project `RUSTUP_TOOLCHAIN` reaches it too. The probe lives next to its one caller, and `backend::active_host_triple_in` goes. - The README placed `--target` before `--`; Cargo arguments follow the first `--`. It now says so, names `CARGO_BUILD_TARGET` for commands without Cargo arguments, and lists `host-tuple`, the unread `[build] target` / `[build] rustc`, and that the probe does not go through `RUSTC_WRAPPER`. The error for an unknown triple now also covers a probe that could not run. Tests: `host-tuple` on the command line and in both env sources, its dedup against the host's name, one probe for several mentions, and the named error when the probe fails; the probe's compiler choice, directory and environment; and a fake compiler whose `host:` line is read, and one that fails. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Signed-off-by: Alis Volat Propriis <alisvolatprop12@proton.me> --------- Signed-off-by: Alis Volat Propriis <alisvolatprop12@proton.me> Co-authored-by: Alis Volat Propriis <alisvolatprop12@proton.me> Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Signed-off-by: juliankang4 <200967947+juliankang4@users.noreply.github.com>
…VIDIA#3) Signed-off-by: juliankang4 <200967947+juliankang4@users.noreply.github.com>
…VIDIA#2) Signed-off-by: Burntgogi <224273819+Burntgogi@users.noreply.github.com>
The branch replays bloomery's patches on NVIDIA/cuda-rust 9df9e7c, so every conflicting file keeps the replayed version: the range `for` unroll and the ordinary-call pointer adaptation as upstream merged them, and the per-target `extra-rustflags` and shared_memory_helper changes rebased onto the upstream code around them. No file outside the conflicts differs from the replayed tree. Signed-off-by: juliankang4 <200967947+juliankang4@users.noreply.github.com>
`OrdinaryFacts` kept one flat map keyed by (value, field path), so `entries(value)` scanned every key in the map. The ordinary plumbing and the poison pass call it for every cast, aggregate, store, load and block argument of a scope, which made each sweep quadratic in the scope's size. Key the facts by SSA value first and by field path second. `entries`, `at_path` and `copy_entries` now touch only the named carrier, and `copy_entries` still forwards a snapshot of the source taken before the copy. The shared-MMA use closure walks the nested map and tests scope membership through a set instead of a linear search. Signed-off-by: juliankang4 <200967947+juliankang4@users.noreply.github.com>
…tions The scheduler, pipeline and shared-MMA stories built their provenance facts for every scope, including every function with no CuTe operation at all. Each analysis seeds every integer value of the scope, so this cost was paid for every ordinary kernel. Return before building the facts when the scope holds no operation of the story's family. Every check in those stories starts from such an operation, so a scope without one was already accepted after the analysis; consumer-only (malformed) stories still reach their checks because their operations belong to the family. Signed-off-by: juliankang4 <200967947+juliankang4@users.noreply.github.com>
Seeding gives every factless integer value an origin that names the value itself (`Unknown`, and `Scalar` for shared MMA). Propagation merges those origins through block arguments and closed local cells, so a value's set grew with the length of every chain that reaches it, and each sweep of the fixed point copied the grown sets: verification time grew with the cube of the scope size. Consumers read a set only through its size, its sole origin, or whether it holds a given semantic origin, and they skip opaque origins when walking a set. Two opaque origins therefore report a merge exactly as any larger number does. `insert_path` now drops a third one, through a `CarrierFact` trait that each family's fact type implements. Path presence and semantic origins are unchanged, so every accept or reject decision and the diagnostic each consumer reports stay the same. Tests cover the cap, the snapshot copy, a scheduler state carried through 48 blocks and cells, and the rejection wording for cursors that merge one to eight unrelated values. Signed-off-by: juliankang4 <200967947+juliankang4@users.noreply.github.com>
Every load and store that forwards provenance through a closed local cell called `local_cell_is_closed`, which walks all uses of the cell and its aliases. A cell with N accesses therefore cost N walks of N uses on every fixed-point sweep, so verification time grew with the square of the number of accesses to one cell. The verifier does not change the IR it reads, so a cell's closedness is fixed for the whole analysis. `LocalCells` walks the alias graph of each `mir.alloca` in the scope once, and the scheduler, pipeline and shared-MMA fact builders, the shared-MMA use closure and the epilogue plumbing check look cells up there. A cell allocated outside the scope is still walked on demand, so every answer is the one the walk gives. The ordinary plumbing and the poison pass also read each incoming edge's successor operands once per block instead of once per block argument. A new test compares the table with the alias walk for closed, aliased and escaped cells, with and without their allocations in the scope. Signed-off-by: juliankang4 <200967947+juliankang4@users.noreply.github.com>
The typed-slot change for aggregate niche payloads and its example came from a branch behind this one, and four checks fail on the result: - The example's manifest put a version on its cuda-core path dependency and its lockfile pinned the 0.4.0 host crates. check-shared-crate-pin.sh rejects a version on a path dependency, and the host crates are 0.5.0. - The lockfile pinned newer registry crates than the other examples. Aligned with scripts/sync-example-locks.sh --align. - The example's main.rs spelled the copyright holder differently from the standard SPDX header, which check-spdx-headers.sh rejects. - enum_slot_map_pointer_niche_with_u128_payload_preserves_align16 expected the u128 bytes as two-i64 filler behind an alignment anchor. The payload now decomposes, so the u128 leaf has its own i128 slot and no anchor is needed; the test checks that layout and keeps the 32-byte, 16-aligned size check. Signed-off-by: juliankang4 <200967947+juliankang4@users.noreply.github.com>
The expected output was recorded with the nightly the fork pinned before. The pinned 1.99.1 compiler shows the similarly named constant as a label on the same snippet instead of a separate note, so trybuild reported a mismatch. Signed-off-by: juliankang4 <200967947+juliankang4@users.noreply.github.com>
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
This moves
bloomeryfrom NVIDIA/cuda-rust1691162ato9df9e7ce, keepsbloomery's own patches on top, and adds the fix in NVIDIA#1648 for the CuTe semantic verifier that upstream added in927d462e(NVIDIA#1598). The branch mergesbloomeryback in, so it merges without conflicts, and its diff againstbloomeryis the upstream move plus the verifier fix.Without the fix, a cold release build of a 358-kernel crate on an RTX 5090 took 220 s at
0143af5c(the parent of927d462e) and did not finish within 1700 s at9df9e7ce. With the first three verifier commits, the same build takes 276 s; the fourth only runs in scopes that use CuTe operations, which this crate has none of.Changes
bloomery's patches on9df9e7ce:feat(cuda-macros): let requires name unsigned integer constants,fix(cuda-macros): name the macro call when a cuda_module finds no kernels,fix(mir-lower): give an aggregate niche payload's leaves typed slotsandcodegen: enforce a project's required llc major on the selected tool (#2)apply unchanged.feat(cargo-oxide): select extra-rustflags by target tripleis rebased onto upstream'scodegen_env.rs: it keeps upstream'scutlass_toolchainimport andapply_profile_pinscall, and its tests follow upstream's newtilelibtests.feat(unroll): recognize range for loopsis not replayed, because upstream merged it as feat(unroll): recognize rangeforloops NVIDIA/cuda-rust#1346 (6d47ae49); the upstream version is kept.fix(mir-importer): adapt shared pointers at ordinary call boundaries (#3): upstream927d462ealready carries the same importer change, so only theexamples-compileworkflow and theshared_memory_helperREADME changes are replayed. The README keeps upstream's paragraph on theSharedPtrprobes.bloomeryrecords it as a parent. Every conflicting file keeps the replayed version, and no other file differs from the replayed tree.Three commits after the verifier fix adapt the replayed patches to the checks at
9df9e7ce:fix(mir-lower): fit enum_niche_payload to the current tree: the example'scuda-corepath dependency drops its version, its lockfile moves to the 0.5.0 host crates and is aligned with the other examples byscripts/sync-example-locks.sh --align, itsmain.rsgets the standard SPDX copyright line, andenum_slot_map_pointer_niche_with_u128_payload_preserves_align16checks the decomposed layout (ani128slot for theu128leaf, the same 32-byte, 16-aligned size).test(cuda-macros): bless the unresolved-constant diagnostic for 1.99.1: the expected trybuild output was recorded with the nightlybloomerypinned before, and 1.99.1 places the same note differently.ci: compile the shared-memory helper calls with the job's llc-23: the replayedexamples-compilestep still namedllc-22, which the job no longer installs since upstream raised the LLVM floor to 23.The verifier fix, in
cuda-oxide/cute-rs-preview/dialect-cute/src/verify/:perf(cute): look up ordinary provenance facts by carrier:OrdinaryFactsis keyed by SSA value and then field path, so lookups no longer scan every fact.perf(cute): skip a family's fact analysis in scopes without its operations: the scheduler, pipeline and shared-MMA stories return early in scopes without their operations.perf(cute): keep at most two opaque origins per provenance fact set: opaque origins no longer pile up along block-argument and local-cell chains.perf(cute): check each local cell's aliases once per analysis: closed local cells are found once per analysis instead of once per access and sweep.The accept and reject rules and the diagnostic each consumer reports stay the same; NVIDIA#1648 lists each consumer of the fact sets and why two opaque origins give the same answer.
Testing
cargo test -p dialect-cute --all-targetson the branch:93 passed; 0 failedin the library tests, integration tests pass.cargo test -p cargo-oxide --all-targetson the branch on macOS:290 passed; 5 failed. The same five tests fail on that machine at9df9e7cewithout these changes (278 passed; 5 failed): three need atarthat accepts--delay-directory-restore, one a non-UTF-8 path, and one a configured target directory.cargo oxide build --arch sm_86 -- --release -j 12, Ubuntu 26.04.1, CUDA 13.3, driver 595.99.02, rustc 1.99.1 (0143af5cused its pinnednightly-2026-10-11).On
37901e4a, RTX 5090 with driver 595.99.02, CUDA 13.3, LLVM 21.1.8 and the pinned 1.99.1 toolchain:just -f cuda-oxide/Justfile checkpasses (exit 0)cargo oxide run vecaddpasses:✓ SUCCESS: All 1024 elements correct!scripts/sync-example-locks.sh --check:OK: 259 example locks agree on one version per crateChecklist
git commit -s)