chore(cuda-oxide): update flake.lock for 1.99.1, scaffold config, and clean up docs

.github/workflows/status-guard.yml

@@ -60,13 +60,12 @@ jobs:
       - name: Verify README lists every workspace member
         run: bash cuda-oxide/scripts/check-crate-inventory.sh
 
-  # The pin is copied into places no build compares: the nested backend pin
-  # (rustup reads it for the backend's own [workspace], so a mismatch there
-  # is a load failure rather than a warning), the scaffold `cargo oxide new`
-  # writes into user projects, the devcontainer's preinstalled toolchain, the
-  # book's quoted [toolchain] blocks, and the dated rustup/cargo commands in
-  # the book and READMEs. #727 added rustfmt to both real files and left both
-  # book quotes behind.
+  # The pin is copied into places no build compares: the scaffold `cargo oxide
+  # new` writes into user projects, the devcontainer's preinstalled toolchain,
+  # the book's quoted [toolchain] blocks, the explicit rustup/cargo commands in
+  # the book and READMEs, and the `rust-overlay` input in `cuda-oxide/flake.lock`.
+  # #727 added rustfmt to both real files and left both book quotes behind;
+  # #1609 bumped the pin to 1.99.1 and left `flake.lock` behind.
   toolchain-pin-parity:
     name: toolchain pin parity
     runs-on: ubuntu-latest
@@ -130,9 +129,9 @@ jobs:
       - name: Clone repository
         uses: actions/checkout@v4
 
-# rustfmt comes from cuda-oxide/rust-toolchain.toml's component list through the
+      # rustfmt comes from cuda-oxide/rust-toolchain.toml's component list through the
       # rustup proxy, as in fmt.yml, so the order checked is the pinned
-      # nightly's. No cargo, no CUDA toolkit, no network beyond the toolchain.
+      # toolchain's. No cargo, no CUDA toolkit, no network beyond the toolchain.
       - name: Verify the book's Rust blocks use rustfmt's import order
         run: bash cuda-oxide/scripts/check-book-import-order.sh
 

.github/workflows/unit-tests.yml

@@ -122,21 +122,6 @@ jobs:
           sub-packages: '["nvcc", "cudart-dev"]'
           non-cuda-sub-packages: '["libcurand-dev"]'
 
-      - name: Install LLVM 21 (llc/opt for the NVPTX backend)
-        if: matrix.package == 'cuda-oxide-codegen'
-        run: |
-          set -euo pipefail
-          source /etc/os-release
-          sudo install -d -m 0755 /etc/apt/keyrings
-          wget -qO- https://apt.llvm.org/llvm-snapshot.gpg.key \
-            | sudo tee /etc/apt/keyrings/apt-llvm-org.asc >/dev/null
-          echo "deb [signed-by=/etc/apt/keyrings/apt-llvm-org.asc] https://apt.llvm.org/${VERSION_CODENAME}/ llvm-toolchain-${VERSION_CODENAME}-21 main" \
-            | sudo tee /etc/apt/sources.list.d/llvm-21.list
-          sudo apt-get update
-          sudo apt-get install -y --no-install-recommends llvm-21
-          llc-21 --version | head -2
-          opt-21 --version | head -2
-
       - name: Run crate unit tests
         env:
           CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}

AGENTS.md

@@ -31,8 +31,9 @@ contribution guide. There is no root `Justfile`. The SIMT recipes live in
 - `cuda-bindings/`, `cuda-core/`, `cuda-core-derive/`, and `cuda-async/` — shared
   host crates in the root workspace. Both kernel stacks depend on these by path.
 - `cuda-oxide/` — the SIMT model. Its `Cargo.toml` explicitly lists workspace
-  members under `crates/`; its `rust-toolchain.toml` pins the nightly needed for
-  `rustc_private`. The book is `cuda-oxide/cuda-oxide-book/`, and
+  members under `crates/`; its `rust-toolchain.toml` pins the stable release
+  (`1.99.1`) and `rustc-dev` / `llvm-tools` components needed for `rustc_private`
+  (unlocked via `RUSTC_BOOTSTRAP=1`). The book is `cuda-oxide/cuda-oxide-book/`, and
   `cuda-oxide/intrinsics/` holds the intrinsic catalog and ABI ledger.
   `cuda-oxide/crates/rustc-codegen-cuda/` declares its own workspace and inherits
   the SIMT toolchain pin. Examples and test fixtures can also declare separate
@@ -84,15 +85,16 @@ nested workspaces.
 ## Toolchains
 
 The root `rust-toolchain.toml` and `cutile-rs/rust-toolchain.toml` carry matching
-stable pins. `cuda-oxide/rust-toolchain.toml` supplies the SIMT nightly, including
-for its nested backend workspace. Do not float a pin to work around a failure.
+stable pins. `cuda-oxide/rust-toolchain.toml` supplies the SIMT stable pin and
+compiler components (`rustc-dev`, `llvm-tools`), including for its nested backend
+workspace. Do not float a pin to work around a failure.
 `cuda-oxide/scripts/check-toolchain-parity.sh` checks the SIMT pin against its
 scaffold, development container, and documentation.
 
 rustup normally resolves the toolchain from the working directory; an explicit
 override can change that. `--manifest-path` does not change the working directory,
 so pointing a root Cargo command at the SIMT manifest does not select the SIMT
-nightly. Change directory first or select the required toolchain explicitly.
+toolchain pin. Change directory first or select the required toolchain explicitly.
 
 Nix environments live in `cuda-oxide/flake.nix` and `cutile-rs/flake.nix`.
 Run `nix develop` from the relevant component.

CONTRIBUTING.md

@@ -142,13 +142,16 @@ Each product has its own workspace and toolchain under `cuda-oxide/` or
 `cutile-rs/`. The setup commands below cover cuda-oxide; the Tile checks have
 their own section further down.
 
-cuda-oxide requires the Rust nightly toolchain with `rustc_private` support.
-See the [cuda-oxide README](cuda-oxide/README.md) for setup instructions.
-
-`cuda-oxide/flake.nix` provides the CUDA 13, LLVM 22, Clang, and pinned nightly
-development environment. From the repository root, run `cd cuda-oxide` and
-then `nix develop` with flakes enabled. Container setup is documented in
-[cuda-oxide/.devcontainer/README.md](cuda-oxide/.devcontainer/README.md).
+cuda-oxide pins its stable Rust toolchain (`1.99.1`) and required `rustc-dev`
+and `llvm-tools` components in `cuda-oxide/rust-toolchain.toml` (with
+`RUSTC_BOOTSTRAP=1` configured via `cuda-oxide/.cargo/config.toml` and
+`cargo-oxide`). See the [cuda-oxide README](cuda-oxide/README.md) for setup
+instructions.
+
+`cuda-oxide/flake.nix` provides the CUDA 13, Clang/libclang, and pinned Rust
+(with LLVM 23 `llvm-tools`) development environment. From the repository root,
+run `cd cuda-oxide` and then `nix develop` with flakes enabled. Container setup
+is documented in [cuda-oxide/.devcontainer/README.md](cuda-oxide/.devcontainer/README.md).
 
 ### Running the cuda-oxide checks
 

cuda-oxide/README.md

@@ -150,7 +150,7 @@ cargo oxide update
 ### Requirements
 
 - **cargo-oxide** — cargo subcommand that drives the build pipeline (`cargo oxide run`, `build`, `sanitize`, `debug`, etc.)
-- **Rust nightly** with `rust-src` and `rustc-dev` and `llvm-tools` components (pinned in `rust-toolchain.toml`)
+- **Rust 1.99.1 (stable)** with `rust-src`, `rustc-dev`, and `llvm-tools` components (pinned in `rust-toolchain.toml`)
 - **CUDA Toolkit** (13.0+, including the cuRAND headers; `libcurand-dev` on Ubuntu). The shared `cuda-bindings` crate loads `libcuda` at run time and needs a CUDA 13.x driver (R580+)
 - **Clang + libclang dev headers** (`clang-21` / `libclang-common-21-dev`) — needed by `bindgen` when building the host `cuda-bindings` crate
 - **Linux** (tested on Ubuntu 24.04)

cuda-oxide/crates/cargo-oxide/src/commands/scaffold.rs

@@ -18,6 +18,10 @@ channel = "1.99.1"
 components = ["rust-src", "rustc-dev", "rust-analyzer", "clippy", "rustfmt", "llvm-tools"]
 "#;
 
+const CARGO_CONFIG_TOML: &str = r#"[env]
+RUSTC_BOOTSTRAP = "1"
+"#;
+
 const SCAFFOLD_GITIGNORE_EXTRA: &[&str] = &[
     "/target/",
     "**/*.bc", // bitcode leftovers not in the clean suffix list
@@ -52,6 +56,7 @@ pub(super) fn scaffold_gitignore() -> String {
 #[derive(Clone, Debug, Eq, PartialEq)]
 pub(super) struct ScaffoldFiles {
     pub(super) cargo_toml: String,
+    pub(super) cargo_config_toml: String,
     rust_toolchain_toml: String,
     pub(super) gitignore: String,
     pub(super) readme: String,
@@ -299,6 +304,7 @@ fn main() -> Result<(), Box<dyn std::error::Error>> {
 pub(super) fn scaffold_files(name: &str, async_mode: bool) -> ScaffoldFiles {
     ScaffoldFiles {
         cargo_toml: scaffold_cargo_toml(name, async_mode),
+        cargo_config_toml: CARGO_CONFIG_TOML.to_string(),
         rust_toolchain_toml: RUST_TOOLCHAIN_TOML.to_string(),
         gitignore: scaffold_gitignore(),
         readme: scaffold_readme(name, async_mode),
@@ -319,10 +325,17 @@ pub fn scaffold_new(name: &str, async_mode: bool) {
         eprintln!("Error creating directory: {}", e);
         std::process::exit(1);
     });
+    let cargo_dir = project_dir.join(".cargo");
+    std::fs::create_dir_all(&cargo_dir).unwrap_or_else(|e| {
+        eprintln!("Error creating directory: {}", e);
+        std::process::exit(1);
+    });
 
     let files = scaffold_files(name, async_mode);
     std::fs::write(project_dir.join("Cargo.toml"), files.cargo_toml)
         .expect("Failed to write Cargo.toml");
+    std::fs::write(cargo_dir.join("config.toml"), files.cargo_config_toml)
+        .expect("Failed to write .cargo/config.toml");
     std::fs::write(
         project_dir.join("rust-toolchain.toml"),
         files.rust_toolchain_toml,

cuda-oxide/crates/cargo-oxide/src/commands/tests.rs

@@ -5741,6 +5741,8 @@ fn scaffold_names_the_canonical_repository() {
                 .contains("git = \"https://github.com/NVIDIA/cuda-rust.git\"")
         );
         assert!(!files.cargo_toml.contains("NVlabs"));
+        assert!(files.cargo_config_toml.contains("[env]"));
+        assert!(files.cargo_config_toml.contains("RUSTC_BOOTSTRAP = \"1\""));
     }
 }
 

cuda-oxide/cuda-oxide-book/compiler/rustc-public.md

@@ -282,7 +282,7 @@ already explicit. The `mir-importer` crate translates this into Pliron IR
 Even with `rustc_public` providing a stable API, the bridge layer
 (`rustc_public_bridge`) is compiled against `librustc_driver.so` and is tied to
 an exact compiler build. Instead of chasing daily `nightly` snapshots,
-`cuda-oxide` pins an official **stable** Rust release via
+cuda-oxide pins an official **stable** Rust release via
 `rust-toolchain.toml`:
 
 ```toml
@@ -291,22 +291,27 @@ channel = "1.99.1"
 components = ["rust-src", "rustc-dev", "rust-analyzer", "clippy", "rustfmt", "llvm-tools"]
 ```
 
+When bumping this pin, run `cuda-oxide/scripts/check-toolchain-parity.sh` to
+verify that `scaffold.rs`, `.devcontainer/devcontainer.json`, the book, and the
+READMEs agree, and run `nix flake update rust-overlay` in `cuda-oxide/` so
+`cuda-oxide/flake.lock` carries the new release manifest.
+
 ```{figure} images/ecosystem-and-bootstrap.svg
 :align: center
 :width: 100%
 
-How custom Rust compiler backends ship in production, and what lives inside an official stable `rustup` release.
+How custom Rust compiler backends ship in production, and what lives inside an official stable rustup release.
 ```
 
-### Why stable `rustc` already has `rustc_private` and `rustc_public`
+### Why stable rustc already has rustc_private and rustc_public
 
-Official stable `rustc` binaries are compiled from the same source tree as
-nightly; nothing is stripped out of the binary. In addition, `rustup` publishes
+Official stable rustc binaries are compiled from the same source tree as
+nightly; nothing is stripped out of the binary. In addition, rustup publishes
 the `rustc-dev` component (`librustc_driver.so`, `librustc_public.rlib`,
 `librustc_public_bridge.rlib`, and `librustc_middle.rlib`) for every stable
 release.
 
-When `cargo-oxide` invokes `cargo` (and inside the workspace via
+When cargo-oxide invokes Cargo (and inside the workspace via
 `cuda-oxide/.cargo/config.toml`), it sets `RUSTC_BOOTSTRAP=1` in the process
 environment. That unlocks `#![feature(rustc_private)]` and `-Z` flags
 (`-Z codegen-backend=...`, `-Z mir-enable-passes=-JumpThreading`,
@@ -318,10 +323,10 @@ stabilization cycle.
 :align: center
 :width: 100%
 
-Pinning an official stable release (`1.99.1`) in `cuda-oxide/rust-toolchain.toml` with `RUSTC_BOOTSTRAP=1` managed by `cargo-oxide`.
+Pinning an official stable release (`1.99.1`) in `cuda-oxide/rust-toolchain.toml` with `RUSTC_BOOTSTRAP=1` managed by cargo-oxide.
 ```
 
-### Upstreaming to `rustc_public` without emergency toolchain bumps
+### Upstreaming to rustc_public without emergency toolchain bumps
 
 When a feature requires a new `rustc_public` query that we upstream to
 `rust-lang/rust` (such as `ArgAttributes` getters in

cuda-oxide/cuda-oxide-book/cute/contracts.md

@@ -8,16 +8,16 @@ with the author at `unsafe` boundaries.
 :align: center
 :width: 100%
 
-Where `ctm-rs` sits: typed carriers select instruction families, one of two compiler profiles consumes them, and the author keeps the schedule.
+Where ctm-rs sits: typed carriers select instruction families, one of two compiler profiles consumes them, and the author keeps the schedule.
 ```
 
 ## What is checked
 
 | Layer | Checks |
 | :--- | :--- |
 | Rust types | layout, swizzle, operand role, element format, MMA shape, pipeline stage count, barrier ids; `TensorMut` over `DisjointSlice`; `TmemGuard` typestate; sealed traits |
-| `cute-layout` proofs | cooperative copy coverage and per-thread contiguity, `ldmatrix` source validity, TMA encodability and swizzle phase alignment |
-| `dialect-cute` verifier | tensor provenance, one pipeline init with complete producer and consumer state, cursor provenance, scheduler grids, MMA placement, epilogue phases, no compiler-only type reaches runtime storage |
+| cute-layout proofs | cooperative copy coverage and per-thread contiguity, `ldmatrix` source validity, TMA encodability and swizzle phase alignment |
+| dialect-cute verifier | tensor provenance, one pipeline init with complete producer and consumer state, cursor provenance, scheduler grids, MMA placement, epilogue phases, no compiler-only type reaches runtime storage |
 | Target gate | `sm_100a` or `sm_120a`; `tcgen05` only on `sm_100a`; MXF4 register MMA only on `sm_120a` |
 | Launch ABI | cubin argument sizes and offsets match the exported MLIR; `#[grid_constant] &TensorMap` is a 128-byte by-value argument aligned to 64 |