diff --git a/.github/workflows/clippy.yml b/.github/workflows/clippy.yml index 0705378b68..24655cfa24 100644 --- a/.github/workflows/clippy.yml +++ b/.github/workflows/clippy.yml @@ -64,6 +64,7 @@ jobs: env: CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }} CARGO_TARGET_DIR: ${{ runner.temp }}/clippy-shared-target + RUSTC_BOOTSTRAP: "1" run: | set -e example_roots=( diff --git a/.github/workflows/fmt.yml b/.github/workflows/fmt.yml index 72b6b91101..40a946953e 100644 --- a/.github/workflows/fmt.yml +++ b/.github/workflows/fmt.yml @@ -44,6 +44,8 @@ jobs: working-directory: cuda-oxide/crates/cuda-macros/tests/device-only - name: Check example and model crate formatting + env: + RUSTC_BOOTSTRAP: "1" run: | set -euo pipefail # `cargo fmt --all` also formats local path dependencies. Running it diff --git a/cuda-oxide/.cargo/config.toml b/cuda-oxide/.cargo/config.toml index 92b7403412..e660191546 100644 --- a/cuda-oxide/.cargo/config.toml +++ b/cuda-oxide/.cargo/config.toml @@ -5,6 +5,9 @@ oxide = "run --package cargo-oxide --" # Environment variables for all builds [env] +# Unlock `-Z` flags and `#![feature(rustc_private)]` on the pinned stable +# release (`rust-toolchain.toml`). +RUSTC_BOOTSTRAP = "1" # libmathdx installation path (for libmathdx-sys crate) # Required only when using --features libmathdx with mathdx_ffi_test. # IMPORTANT: Replace the placeholder below with your libmathdx install dir, diff --git a/cuda-oxide/.devcontainer/README.md b/cuda-oxide/.devcontainer/README.md index 01bcde0c46..12e9b7f612 100644 --- a/cuda-oxide/.devcontainer/README.md +++ b/cuda-oxide/.devcontainer/README.md @@ -6,7 +6,7 @@ This dev container provides the toolchain expected by cuda-oxide: - CUDA Toolkit 13.0 - LLVM 21 with NVPTX support - Clang 21 resource headers for `bindgen` -- Rust `nightly-2026-10-11` with `rust-src`, `rustc-dev`, `rust-analyzer`, +- Rust `1.99.1` with `rust-src`, `rustc-dev`, `rust-analyzer`, `rustfmt` and `clippy`. `llvm-tools` is not preinstalled in the image; rustup adds it on your first cargo command under `cuda-oxide/`, because `cuda-oxide/rust-toolchain.toml` names it. The image ships LLVM 21 and diff --git a/cuda-oxide/.devcontainer/devcontainer.json b/cuda-oxide/.devcontainer/devcontainer.json index 06fa7b4ba7..e77f29ca6e 100644 --- a/cuda-oxide/.devcontainer/devcontainer.json +++ b/cuda-oxide/.devcontainer/devcontainer.json @@ -17,7 +17,7 @@ "upgradePackages": false }, "ghcr.io/devcontainers/features/rust:1": { - "version": "nightly-2026-10-11", + "version": "1.99.1", "profile": "minimal", "components": "rust-src,rustc-dev,rust-analyzer,rustfmt,clippy" } diff --git a/cuda-oxide/Justfile b/cuda-oxide/Justfile index 9df0c0f1bd..2690e9f273 100644 --- a/cuda-oxide/Justfile +++ b/cuda-oxide/Justfile @@ -7,6 +7,10 @@ set windows-shell := ["powershell.exe", "-NoLogo", "-Command"] # Use Bash on Unix set shell := ["bash", "-c"] +# Unlock unstable Cargo (`cargo-features = ["profile-rustflags"]`) and rustc +# (`-Z`, `#![feature(rustc_private)]`) gates on the pinned stable toolchain. +export RUSTC_BOOTSTRAP := "1" + # Format all Rust code (SIMT workspace, codegen, examples) fmt: cargo oxide fmt diff --git a/cuda-oxide/README.md b/cuda-oxide/README.md index ed157bd6c0..79f8349b6b 100644 --- a/cuda-oxide/README.md +++ b/cuda-oxide/README.md @@ -161,17 +161,17 @@ cargo oxide update From `cuda-oxide/`, `cargo oxide` works out of the box via a workspace alias. -For use outside the repo (your own projects), install it with the pinned nightly toolchain: +For use outside the repo (your own projects), install it with the pinned toolchain: ```bash -cargo +nightly-2026-10-11 install --git https://github.com/NVIDIA/cuda-rust.git cargo-oxide +cargo +1.99.1 install --git https://github.com/NVIDIA/cuda-rust.git cargo-oxide ``` On first run, `cargo-oxide` will automatically fetch and build the codegen backend. #### Nix (alternative) -If you have Nix with flakes enabled, `nix develop` from `cuda-oxide/` gives you a reproducible shell with CUDA 13, LLVM 22, Clang, and the pinned Rust nightly — no manual apt installs. The shellHook auto-discovers host NVIDIA drivers on NixOS and non-NixOS systems. +If you have Nix with flakes enabled, `nix develop` from `cuda-oxide/` gives you a reproducible shell with CUDA 13, LLVM 22, Clang, and the pinned Rust toolchain - no manual apt installs. The shellHook auto-discovers host NVIDIA drivers on NixOS and non-NixOS systems. ```bash # From the repository root: @@ -187,8 +187,8 @@ nix run 'github:NVIDIA/cuda-rust?dir=cuda-oxide#new' -- my-project ```bash # Toolchain installed automatically via rust-toolchain.toml # Manual install if needed: -rustup toolchain install nightly-2026-10-11 -rustup component add rust-src rustc-dev rust-analyzer clippy rustfmt llvm-tools --toolchain nightly-2026-10-11 +rustup toolchain install 1.99.1 +rustup component add rust-src rustc-dev rust-analyzer clippy rustfmt llvm-tools --toolchain 1.99.1 ``` #### CUDA diff --git a/cuda-oxide/crates/cargo-oxide/README.md b/cuda-oxide/crates/cargo-oxide/README.md index 71a95180a2..f9158e3c39 100644 --- a/cuda-oxide/crates/cargo-oxide/README.md +++ b/cuda-oxide/crates/cargo-oxide/README.md @@ -10,10 +10,10 @@ Replaces the previous `xtask` pattern with a proper cargo subcommand that works **External users**: -Install with the project's pinned nightly toolchain: +Install with the project's pinned toolchain: ```bash -cargo +nightly-2026-10-11 install --git https://github.com/NVIDIA/cuda-rust.git cargo-oxide +cargo +1.99.1 install --git https://github.com/NVIDIA/cuda-rust.git cargo-oxide ``` On first run, `cargo-oxide` will automatically fetch and build the codegen backend if it's not already available. diff --git a/cuda-oxide/crates/cargo-oxide/src/backend.rs b/cuda-oxide/crates/cargo-oxide/src/backend.rs index 2e5ce0d5ff..37d63ebea1 100644 --- a/cuda-oxide/crates/cargo-oxide/src/backend.rs +++ b/cuda-oxide/crates/cargo-oxide/src/backend.rs @@ -1056,6 +1056,7 @@ fn backend_build_command_in( cmd.env("CARGO_ENCODED_RUSTFLAGS", ""); cmd.env_remove("RUSTFLAGS"); cmd.env_remove("CARGO_BUILD_RUSTFLAGS"); + cmd.env("RUSTC_BOOTSTRAP", "1"); if let Some(path) = rustc_lib_path { cmd.env("LIBRARY_PATH", build_library_path(path)); diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/codegen_env.rs b/cuda-oxide/crates/cargo-oxide/src/commands/codegen_env.rs index 86fb90a163..1442a2ce61 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/codegen_env.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/codegen_env.rs @@ -466,6 +466,7 @@ fn apply_codegen_rustflags( append_full_debug_rustflags(&mut encoded, cmd, inherited_debug.as_deref()); cmd.env("CARGO_ENCODED_RUSTFLAGS", encoded) + .env("RUSTC_BOOTSTRAP", "1") .env_remove("RUSTFLAGS"); apply_profile_pins(cmd, profile); } @@ -553,6 +554,7 @@ pub fn has_configured_arch(ctx: &Context, cli_arch: Option<&str>) -> bool { } pub(super) fn apply_config_env(cmd: &mut Command, ctx: &Context) { + cmd.env("RUSTC_BOOTSTRAP", "1"); for (key, value) in &ctx.config.env { if matches!(key.as_str(), "RUSTFLAGS" | "CARGO_ENCODED_RUSTFLAGS") { continue; diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/fmt.rs b/cuda-oxide/crates/cargo-oxide/src/commands/fmt.rs index 79081a32ad..00bb589f16 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/fmt.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/fmt.rs @@ -144,6 +144,7 @@ fn run_cargo_fmt_manifest(manifest: &Path, check: bool) -> bool { } fn run_fmt_command(mut cmd: Command) -> bool { + cmd.env("RUSTC_BOOTSTRAP", "1"); match cmd.status() { Ok(status) => status.success(), Err(e) => { diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/host_cargo.rs b/cuda-oxide/crates/cargo-oxide/src/commands/host_cargo.rs index 36b1f36724..179ad7bb6f 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/host_cargo.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/host_cargo.rs @@ -367,6 +367,7 @@ fn cargo_package_manifest_matches(package: &serde_json::Value, manifest_path: &P fn cargo_metadata(manifest_dir: &Path) -> Result { let output = Command::new("cargo") + .env("RUSTC_BOOTSTRAP", "1") .args(["metadata", "--format-version=1", "--no-deps"]) .current_dir(manifest_dir) .output() diff --git a/cuda-oxide/crates/cargo-oxide/src/commands/scaffold.rs b/cuda-oxide/crates/cargo-oxide/src/commands/scaffold.rs index db77c02d29..3a317d78cf 100644 --- a/cuda-oxide/crates/cargo-oxide/src/commands/scaffold.rs +++ b/cuda-oxide/crates/cargo-oxide/src/commands/scaffold.rs @@ -14,7 +14,7 @@ use super::*; const GIT_REPO: &str = "https://github.com/NVIDIA/cuda-rust.git"; const RUST_TOOLCHAIN_TOML: &str = r#"[toolchain] -channel = "nightly-2026-10-11" +channel = "1.99.1" components = ["rust-src", "rustc-dev", "rust-analyzer", "clippy", "rustfmt", "llvm-tools"] "#; diff --git a/cuda-oxide/crates/cuda-device/src/atomic.rs b/cuda-oxide/crates/cuda-device/src/atomic.rs index 7433fec228..fa616afd12 100644 --- a/cuda-oxide/crates/cuda-device/src/atomic.rs +++ b/cuda-oxide/crates/cuda-device/src/atomic.rs @@ -494,7 +494,7 @@ define_float_atomic! { /// /// `swap` lowers to hardware `atom.exch.b32`; this API does not expose /// `compare_exchange` for floats. `fetch_add` uses LLVM `atomicrmw fadd`. - /// With the pinned `nightly-2026-10-11` toolchain's LLVM 23 NVPTX defaults, + /// With the pinned `1.99.1` toolchain's LLVM 23 NVPTX defaults, /// global and generic `f32` adds use a compare-and-swap loop; shared `f32` /// and `f64` adds still use native instructions on supported targets. /// diff --git a/cuda-oxide/crates/cuda-device/tests/compile_fail/swizzle_non_involutive.stderr b/cuda-oxide/crates/cuda-device/tests/compile_fail/swizzle_non_involutive.stderr index b5d29acc9d..599474d478 100644 --- a/cuda-oxide/crates/cuda-device/tests/compile_fail/swizzle_non_involutive.stderr +++ b/cuda-oxide/crates/cuda-device/tests/compile_fail/swizzle_non_involutive.stderr @@ -6,8 +6,6 @@ error[E0080]: evaluation panicked: Swizzle requires M + B <= S, otherwi | ::: src/swizzle.rs | - | impl Swizzle { - | /// Compile-time check that the swizzle is an involution. | const INVOLUTION: () = assert!( | ____________________________- | | M + B <= S, diff --git a/cuda-oxide/crates/cuda-device/tests/compile_fail/thread_index_runtime_mismatch.stderr b/cuda-oxide/crates/cuda-device/tests/compile_fail/thread_index_runtime_mismatch.stderr index fb4421dbec..4ca8d8067f 100644 --- a/cuda-oxide/crates/cuda-device/tests/compile_fail/thread_index_runtime_mismatch.stderr +++ b/cuda-oxide/crates/cuda-device/tests/compile_fail/thread_index_runtime_mismatch.stderr @@ -13,17 +13,3 @@ note: method defined here | | pub fn get_mut<'kernel>(&mut self, idx: ThreadIndex<'kernel, IndexSpace>) -> Option<&mut T> { | ^^^^^^^ -note: the `get_mut` call is resolved to the method in `cuda_device::DisjointSlice<'a, T, IndexSpace>`, shadowing the method of the same name on the inherent impl for `cuda_device::DisjointSlice<'a, T, cuda_device::Runtime2DIndex>` - --> tests/compile_fail/thread_index_runtime_mismatch.rs:18:17 - | -18 | let _ = out.get_mut(idx); - | ^^^^^^^ refers to `cuda_device::DisjointSlice::<'a, T, IndexSpace>::get_mut` - = note: additionally, there is 1 other available method that isn't in scope -help: you might have meant to call one of the other methods; you can use the fully-qualified path to call one of them explicitly - | -18 - let _ = out.get_mut(idx); -18 + let _ = DisjointSlice::get_mut(&mut out, idx); - | -18 - let _ = out.get_mut(idx); -18 + let _ = <_ as std::slice::SliceIndex<_>>::get_mut(out, idx); - | diff --git a/cuda-oxide/crates/cuda-device/tests/compile_fail/thread_index_stride_mismatch.stderr b/cuda-oxide/crates/cuda-device/tests/compile_fail/thread_index_stride_mismatch.stderr index 96a0802bce..7eed21fb06 100644 --- a/cuda-oxide/crates/cuda-device/tests/compile_fail/thread_index_stride_mismatch.stderr +++ b/cuda-oxide/crates/cuda-device/tests/compile_fail/thread_index_stride_mismatch.stderr @@ -13,17 +13,3 @@ note: method defined here | | pub fn get_mut<'kernel>(&mut self, idx: ThreadIndex<'kernel, IndexSpace>) -> Option<&mut T> { | ^^^^^^^ -note: the `get_mut` call is resolved to the method in `cuda_device::DisjointSlice<'a, T, IndexSpace>`, shadowing the method of the same name on the inherent impl for `cuda_device::DisjointSlice<'a, T, cuda_device::Runtime2DIndex>` - --> tests/compile_fail/thread_index_stride_mismatch.rs:11:17 - | -11 | let _ = out.get_mut(idx); - | ^^^^^^^ refers to `cuda_device::DisjointSlice::<'a, T, IndexSpace>::get_mut` - = note: additionally, there is 1 other available method that isn't in scope -help: you might have meant to call one of the other methods; you can use the fully-qualified path to call one of them explicitly - | -11 - let _ = out.get_mut(idx); -11 + let _ = DisjointSlice::get_mut(&mut out, idx); - | -11 - let _ = out.get_mut(idx); -11 + let _ = <_ as std::slice::SliceIndex<_>>::get_mut(out, idx); - | diff --git a/cuda-oxide/crates/cuda-host/src/entry_registry.rs b/cuda-oxide/crates/cuda-host/src/entry_registry.rs index 657da3159a..b12885c3af 100644 --- a/cuda-oxide/crates/cuda-host/src/entry_registry.rs +++ b/cuda-oxide/crates/cuda-host/src/entry_registry.rs @@ -195,7 +195,7 @@ fn type_id_divergence_message( format!( "generic kernel PTX entry `{requested}` was not found, but the loaded module DOES \ contain `{base}` under a different type-identity hash: {found}. The host \ - (core::intrinsics::reflection::type_id) and the device backend (tcx.type_id_hash) disagreed on \ + (core::intrinsics::type_id) and the device backend (tcx.type_id_hash) disagreed on \ the kernel's `_TID_` name, which means the TypeId-to-u128 extraction in \ cuda_host::type_id is no longer value-correct on this toolchain. Host and device \ code must be built by the same pinned nightly; re-verify the TypeId contract in \ @@ -329,10 +329,7 @@ mod tests { "{message}" ); // The message must identify the two hash producers and the remedy. - assert!( - message.contains("core::intrinsics::reflection::type_id"), - "{message}" - ); + assert!(message.contains("core::intrinsics::type_id"), "{message}"); assert!(message.contains("tcx.type_id_hash"), "{message}"); assert!(message.contains("same pinned nightly"), "{message}"); assert!(message.contains("compute_kernel_export_name"), "{message}"); diff --git a/cuda-oxide/crates/cuda-host/src/launch.rs b/cuda-oxide/crates/cuda-host/src/launch.rs index c3a2266731..c6033b1dc1 100644 --- a/cuda-oxide/crates/cuda-host/src/launch.rs +++ b/cuda-oxide/crates/cuda-host/src/launch.rs @@ -164,7 +164,7 @@ pub trait CudaKernel { /// The bound on `ptx_name` is intentionally just whatever the kernel /// itself declared — typically `Copy` on the value-passed generics. No /// `'static` is added: the helper `cuda_host::type_id_u128` uses -/// `core::intrinsics::reflection::type_id`, which is bound `T: ?Sized` rather than +/// `core::intrinsics::type_id`, which is bound `T: ?Sized` rather than /// `T: 'static`, so closures that borrow non-`'static` data still go /// through the typed launch path (just like they did under the legacy /// `type_name`-based scheme). diff --git a/cuda-oxide/crates/cuda-host/src/lib.rs b/cuda-oxide/crates/cuda-host/src/lib.rs index 1fa90924a7..5e6a866a7a 100644 --- a/cuda-oxide/crates/cuda-host/src/lib.rs +++ b/cuda-oxide/crates/cuda-host/src/lib.rs @@ -3,7 +3,7 @@ * SPDX-License-Identifier: Apache-2.0 */ -// Required for `type_id_u128`'s use of `core::intrinsics::reflection::type_id`. The +// Required for `type_id_u128`'s use of `core::intrinsics::type_id`. The // intrinsic is the only way to obtain the same 128-bit hash the backend // uses while keeping the bound at `T: ?Sized` (stable `TypeId::of` would // force `T: 'static` on every kernel marker — see `type_id.rs` for why diff --git a/cuda-oxide/crates/cuda-host/src/type_id.rs b/cuda-oxide/crates/cuda-host/src/type_id.rs index e9aa917b55..6d480a4bf4 100644 --- a/cuda-oxide/crates/cuda-host/src/type_id.rs +++ b/cuda-oxide/crates/cuda-host/src/type_id.rs @@ -10,7 +10,7 @@ //! API would force a `T: 'static` bound on the kernel marker, which would in //! turn reject perfectly valid non-`'static` borrowing closures (e.g. a kernel //! launcher capturing `&[f32]` from a stack frame the caller keeps alive -//! across the launch). The `core::intrinsics::reflection::type_id` form has bound +//! across the launch). The `core::intrinsics::type_id` form has bound //! `T: ?Sized` — i.e. no `'static` requirement — and produces the exact same //! 128-bit value that `tcx.type_id_hash` does for that type, because both go //! through the same `erase_and_anonymize_regions` + stable-hash pipeline. @@ -20,10 +20,10 @@ //! every monomorphized type and const argument, so one hash covers the complete //! specialization without inventing a parallel encoding for const values. //! -//! Framing note for future contributors: `core::intrinsics::reflection::type_id` is an +//! Framing note for future contributors: `core::intrinsics::type_id` is an //! internal API and requires `#![feature(core_intrinsics)]` on the owning //! crate. cuda-oxide already ships against `rustc_private` and pins a -//! nightly toolchain, so this is inside our existing risk surface — but the +//! toolchain, so this is inside our existing risk surface, but the //! helper cannot be lifted into a stable-feeling utility crate without //! re-introducing the feature gate there. @@ -54,7 +54,7 @@ static GENERIC_KERNEL_NAMES: OnceLock /// caller (the borrow must outlive `stream.synchronize()`). #[inline] pub fn type_id_u128() -> u128 { - let id = const { core::intrinsics::reflection::type_id::() }; + let id = const { core::intrinsics::type_id::() }; unsafe { core::mem::transmute::(id) } } diff --git a/cuda-oxide/crates/cuda-intrinsics-gen/src/probe.rs b/cuda-oxide/crates/cuda-intrinsics-gen/src/probe.rs index a4e5520c43..eb1f387362 100644 --- a/cuda-oxide/crates/cuda-intrinsics-gen/src/probe.rs +++ b/cuda-oxide/crates/cuda-intrinsics-gen/src/probe.rs @@ -1799,6 +1799,22 @@ fn parse_ptxas_version(stdout: &str, stderr: &str) -> Result { Ok(version.to_owned()) } +fn llvm_major_minor_series(version: &str) -> Option<&str> { + let rest = version.strip_prefix("LLVM version ")?; + let mut parts = rest.split(['.', '-']); + let major = parts.next()?; + let minor = parts.next()?; + if major.is_empty() + || !major.bytes().all(|b| b.is_ascii_digit()) + || minor.is_empty() + || !minor.bytes().all(|b| b.is_ascii_digit()) + { + return None; + } + let prefix_len = "LLVM version ".len() + major.len() + 1 + minor.len(); + Some(&version[..prefix_len]) +} + fn validate_backend_identity( expected_version: &str, expected_commit: &RustcCommit, @@ -1807,8 +1823,12 @@ fn validate_backend_identity( let ProbeBackendIdentity::SelectedEvidence(actual) = actual else { return Ok(()); }; + let same_version_or_series = actual.tool.version == expected_version + || (llvm_major_minor_series(&actual.tool.version).is_some() + && llvm_major_minor_series(&actual.tool.version) + == llvm_major_minor_series(expected_version)); ensure!( - actual.tool.version == expected_version, + same_version_or_series, "rust-toolchain llc version mismatch: selected evidence records {expected_version:?}, found {:?}; use an explicit `--llc` only for a comparison probe", actual.tool.version ); diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_alias_fakes_row_width.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_alias_fakes_row_width.stderr index 6a068a9dd8..a4cedec322 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_alias_fakes_row_width.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_alias_fakes_row_width.stderr @@ -12,4 +12,5 @@ note: required by a bound in `LoadedModule::alias_fakes_row_width` | ^^^^^^^^^^^^^^ required by this bound in `LoadedModule::alias_fakes_row_width` ... 22 | pub fn alias_fakes_row_width(mut out: DisjointSlice) { + | --------------------- required by a bound in this associated function = note: this error originates in the attribute macro `cuda_module` (in Nightly builds, run with -Z macro-backtrace for more info) diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_alias_hides_row_width.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_alias_hides_row_width.stderr index 1e2dd6faa7..4bc77dcf68 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_alias_hides_row_width.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_alias_hides_row_width.stderr @@ -12,4 +12,5 @@ note: required by a bound in `LoadedModule::alias_hides_row_width` | ^^^^^^^^^^^^^^ required by this bound in `LoadedModule::alias_hides_row_width` ... 22 | pub fn alias_hides_row_width(mut out: DisjointSlice) { + | --------------------- required by a bound in this associated function = note: this error originates in the attribute macro `cuda_module` (in Nightly builds, run with -Z macro-backtrace for more info) diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_fake_disjoint_slice.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_fake_disjoint_slice.stderr index a0d8471d81..530af8ce95 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_fake_disjoint_slice.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_fake_disjoint_slice.stderr @@ -26,6 +26,7 @@ note: required by a bound in `LoadedModule::prepare_lookalike` | ^^^^^^^^^^^^^^ required by this bound in `LoadedModule::prepare_lookalike` ... 21 | pub fn lookalike(mut out: DisjointSlice) { + | --------- required by a bound in this associated function = note: this error originates in the attribute macro `cuda_module` (in Nightly builds, run with -Z macro-backtrace for more info) error[E0277]: the trait bound `for<'__cuda_oxide_disjoint> DisjointSlice<'__cuda_oxide_disjoint, u32>: cuda_device::disjoint::__LaunchContractDisjointSliceAbi` is not satisfied @@ -56,4 +57,5 @@ note: required by a bound in `LoadedModule::prepare_lookalike` | ^^^^^^^^^^^^^^ required by this bound in `LoadedModule::prepare_lookalike` ... 21 | pub fn lookalike(mut out: DisjointSlice) { + | --------- required by a bound in this associated function = note: this error originates in the attribute macro `cuda_module` (in Nightly builds, run with -Z macro-backtrace for more info) diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_misleading_index_alias.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_misleading_index_alias.stderr index 3208581be6..50bc23cd6f 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_misleading_index_alias.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_misleading_index_alias.stderr @@ -22,4 +22,5 @@ note: required by a bound in `LoadedModule::prepare_misleading_name` | ^^^^^^^^^^^^^^ required by this bound in `LoadedModule::prepare_misleading_name` ... 13 | pub fn misleading_name(mut out: DisjointSlice) { + | --------------- required by a bound in this associated function = note: this error originates in the attribute macro `cuda_module` (in Nightly builds, run with -Z macro-backtrace for more info) diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_reordered_disjoint_alias.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_reordered_disjoint_alias.stderr index 851917094c..6509665d5d 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_reordered_disjoint_alias.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/launch_contract_reordered_disjoint_alias.stderr @@ -22,6 +22,7 @@ note: required by a bound in `LoadedModule::prepare_reordered` | ^^^^^^^^^^^^^^ required by this bound in `LoadedModule::prepare_reordered` ... 17 | pub fn reordered(mut out: DisjointSlice) { + | --------- required by a bound in this associated function = note: this error originates in the attribute macro `cuda_module` (in Nightly builds, run with -Z macro-backtrace for more info) error[E0277]: the trait bound `for<'__cuda_oxide_disjoint> cuda_device::DisjointSlice<'__cuda_oxide_disjoint, u64>: cuda_device::disjoint::__LaunchContractDisjointSliceAbi` is not satisfied @@ -48,4 +49,5 @@ note: required by a bound in `LoadedModule::prepare_reordered` | ^^^^^^^^^^^^^^ required by this bound in `LoadedModule::prepare_reordered` ... 17 | pub fn reordered(mut out: DisjointSlice) { + | --------- required by a bound in this associated function = note: this error originates in the attribute macro `cuda_module` (in Nightly builds, run with -Z macro-backtrace for more info) diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_invalid_unroll.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_invalid_unroll.stderr index fe4d9eb24e..d79f0da8ea 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_invalid_unroll.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_invalid_unroll.stderr @@ -1,7 +1,6 @@ error[E0080]: evaluation panicked: partial unroll factor must be in 2..=1024, or 0 for full unrolling --> $WORKSPACE/crates/cuda-device/src/thread.rs | - | pub fn __unroll_config() { | const { validate_unroll_factor(FACTOR) } | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ evaluation of `cuda_device::__unroll_config::<1>::{constant#0}` failed inside this call | diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_unresolved.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_unresolved.stderr index 970bdf2d2e..6cf8574256 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_unresolved.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_unresolved.stderr @@ -1,14 +1,3 @@ -warning: `feature(generic_const_exprs)` is not supported with the next-generation trait solver - --> tests/compile_fail/policy_config_unresolved.rs:4:12 - | -4 | #![feature(generic_const_exprs)] - | ^^^^^^^^^^^^^^^^^^^ - | - = note: `-Znext-solver=globally` is currently enabled by default for testing - = note: reverted the setting to `-Znext-solver=coherence` for this crate - = note: the currently stable trait solver will be used for this crate - = note: see issues #160895 for more information - error[E0599]: no associated function or constant named `MAX_THREADS` found for type parameter `P` in the current scope --> tests/compile_fail/policy_config_unresolved.rs:10:20 | diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_zero_threads.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_zero_threads.stderr index b6d60f72e7..9d71df7e96 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_zero_threads.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_config_zero_threads.stderr @@ -1,7 +1,6 @@ error[E0080]: evaluation panicked: launch_bounds maximum threads must be greater than zero --> $WORKSPACE/crates/cuda-device/src/thread.rs | - | pub fn __launch_bounds_config() { | const { validate_launch_bounds(MAX_THREADS) } | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ evaluation of `cuda_device::__launch_bounds_config::<0, 0>::{constant#0}` failed inside this call | diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_exact_block_too_large.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_exact_block_too_large.stderr index e65b8597b7..abdd18341c 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_exact_block_too_large.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_exact_block_too_large.stderr @@ -1,14 +1,3 @@ -warning: `feature(generic_const_exprs)` is not supported with the next-generation trait solver - --> tests/compile_fail/policy_contract_exact_block_too_large.rs:4:12 - | -4 | #![feature(generic_const_exprs)] - | ^^^^^^^^^^^^^^^^^^^ - | - = note: `-Znext-solver=globally` is currently enabled by default for testing - = note: reverted the setting to `-Znext-solver=coherence` for this crate - = note: the currently stable trait solver will be used for this crate - = note: see issues #160895 for more information - error[E0080]: evaluation panicked: launch_contract exact block exceeds launch_bounds maximum threads --> tests/compile_fail/policy_contract_exact_block_too_large.rs:20:1 | diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_large_3d_block.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_large_3d_block.stderr index 0b2d131856..74e3e4d53a 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_large_3d_block.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_large_3d_block.stderr @@ -1,14 +1,3 @@ -warning: `feature(generic_const_exprs)` is not supported with the next-generation trait solver - --> tests/compile_fail/policy_contract_large_3d_block.rs:4:12 - | -4 | #![feature(generic_const_exprs)] - | ^^^^^^^^^^^^^^^^^^^ - | - = note: `-Znext-solver=globally` is currently enabled by default for testing - = note: reverted the setting to `-Znext-solver=coherence` for this crate - = note: the currently stable trait solver will be used for this crate - = note: see issues #160895 for more information - error[E0080]: evaluation panicked: launch_contract exact block exceeds launch_bounds maximum threads --> tests/compile_fail/policy_contract_large_3d_block.rs:20:1 | diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_wrong_brand.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_wrong_brand.stderr index 4769fb9425..b08b165228 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_wrong_brand.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_wrong_brand.stderr @@ -1,14 +1,3 @@ -warning: `feature(generic_const_exprs)` is not supported with the next-generation trait solver - --> tests/compile_fail/policy_contract_wrong_brand.rs:4:12 - | -4 | #![feature(generic_const_exprs)] - | ^^^^^^^^^^^^^^^^^^^ - | - = note: `-Znext-solver=globally` is currently enabled by default for testing - = note: reverted the setting to `-Znext-solver=coherence` for this crate - = note: the currently stable trait solver will be used for this crate - = note: see issues #160895 for more information - error[E0308]: mismatched types --> tests/compile_fail/policy_contract_wrong_brand.rs:41:16 | diff --git a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_zero_threads.stderr b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_zero_threads.stderr index 7b867238bb..a46d8be4a4 100644 --- a/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_zero_threads.stderr +++ b/cuda-oxide/crates/cuda-macros/tests/compile_fail/policy_contract_zero_threads.stderr @@ -1,14 +1,3 @@ -warning: `feature(generic_const_exprs)` is not supported with the next-generation trait solver - --> tests/compile_fail/policy_contract_zero_threads.rs:4:12 - | -4 | #![feature(generic_const_exprs)] - | ^^^^^^^^^^^^^^^^^^^ - | - = note: `-Znext-solver=globally` is currently enabled by default for testing - = note: reverted the setting to `-Znext-solver=coherence` for this crate - = note: the currently stable trait solver will be used for this crate - = note: see issues #160895 for more information - error[E0080]: evaluation panicked: launch_bounds maximum threads must be greater than zero --> tests/compile_fail/policy_contract_zero_threads.rs:20:1 | diff --git a/cuda-oxide/crates/mir-importer/src/lib.rs b/cuda-oxide/crates/mir-importer/src/lib.rs index 30e2b7f61b..796f34e5cc 100644 --- a/cuda-oxide/crates/mir-importer/src/lib.rs +++ b/cuda-oxide/crates/mir-importer/src/lib.rs @@ -86,6 +86,7 @@ extern crate rustc_middle; extern crate rustc_public; extern crate rustc_public_bridge; extern crate rustc_span; +extern crate rustc_target; /// Value used for every rustc `Operand::RuntimeChecks` query in device MIR. /// diff --git a/cuda-oxide/crates/mir-importer/src/translator/body.rs b/cuda-oxide/crates/mir-importer/src/translator/body.rs index cd49e86c91..2af4b2c883 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/body.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/body.rs @@ -2154,8 +2154,9 @@ pub fn translate_body( for arg_idx in 0..num_args { let local = mir::Local::from(arg_idx + 1); let source_ty = &body.locals()[local].ty; - let Some(validity) = facts::reference_param_validity(source_ty, &fn_abi, arg_idx) - .map_err(|error| input_error_noloc!(TranslationErr::invalid_op(error)))? + let Some(validity) = + facts::reference_param_validity(instance, source_ty, &fn_abi, arg_idx) + .map_err(|error| input_error_noloc!(TranslationErr::invalid_op(error)))? else { continue; }; diff --git a/cuda-oxide/crates/mir-importer/src/translator/facts.rs b/cuda-oxide/crates/mir-importer/src/translator/facts.rs index c9542c38ff..a6b24ac367 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/facts.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/facts.rs @@ -505,6 +505,7 @@ pub(crate) struct ReferenceParamValidity { /// `DisjointSlice` produce no fact. Unsupported reference pass modes and source /// argument mismatches are errors rather than guesses about the ABI mapping. pub(crate) fn reference_param_validity( + instance: &rustc_public::mir::mono::Instance, ty: &rustc_public::ty::Ty, fn_abi: &rustc_public::abi::FnAbi, source_index: usize, @@ -542,15 +543,45 @@ pub(crate) fn reference_param_validity( return Ok(None); } - let attrs = match (&argument.mode, is_slice) { - (PassMode::Direct(attrs), false) | (PassMode::Pair(attrs, _), true) => attrs, + match (&argument.mode, is_slice) { + (PassMode::Direct(_), false) | (PassMode::Pair(_, _), true) => {} _ => { return Err(format!( "kernel reference argument {source_index} has an unsupported rustc FnAbi pass mode" )); } - }; - let flags = attrs.regular(); + } + + let (no_alias, no_free, pointee_size_bytes) = rustc_middle::ty::tls::with(|tcx| { + use rustc_middle::ty::{List, TypingEnv}; + use rustc_target::callconv::{ArgAttribute, PassMode as InternalPassMode}; + + let internal_instance = rustc_public::rustc_internal::internal(tcx, *instance); + let query_input = + TypingEnv::fully_monomorphized().as_query_input((internal_instance, List::empty())); + let internal_fn_abi = tcx + .fn_abi_of_instance(query_input) + .map_err(|err| format!("could not query internal FnAbi: {err:?}"))?; + let internal_arg = internal_fn_abi.args.get(source_index).ok_or_else(|| { + format!("kernel source argument index {source_index} is absent from internal FnAbi") + })?; + let attrs = match (&internal_arg.mode, is_slice) { + (InternalPassMode::Direct(attrs), false) | (InternalPassMode::Pair(attrs, _), true) => { + attrs + } + _ => { + return Err(format!( + "kernel reference argument {source_index} has an unsupported rustc FnAbi pass mode" + )); + } + }; + Ok(( + attrs.regular.contains(ArgAttribute::NoAlias), + attrs.regular.contains(ArgAttribute::NoFree), + attrs.pointee_size.bytes(), + )) + })?; + // Do not consume rustc's `read_only`: NVPTX (including libNVVM) can lower // volatile reads through a readonly + noalias kernel parameter to cached // non-volatile loads. Let the optimizer infer read-only access from the @@ -560,14 +591,10 @@ pub(crate) fn reference_param_validity( // dereferenceability as lasting for the whole call. rustc's NoFree is the // proof needed by all supported consumers; producer LLVM version alone is // insufficient. Zero covers both ZST pointees and unknown slice lengths. - let dereferenceable_bytes = if flags.no_free { - attrs.pointee_size().bytes() as u64 - } else { - 0 - }; + let dereferenceable_bytes = if no_free { pointee_size_bytes } else { 0 }; Ok(Some(ReferenceParamValidity { pointee_alignment: shape.abi_align, - no_alias: flags.no_alias, + no_alias, dereferenceable_bytes, })) } @@ -1055,7 +1082,7 @@ pub fn reference_validity( .iter() .enumerate() .map(|(index, ty)| { - reference_param_validity(ty, &fn_abi, index) + reference_param_validity(&instance, ty, &fn_abi, index) .expect("source parameter matches ABI") }) .collect::>(); @@ -1063,19 +1090,27 @@ pub fn reference_validity( // Reusing a neighboring argument after dropping an ignored // ZST or flattening a slice must not mint reference facts. let mut errors = vec![ - reference_param_validity(&source_types[1], &fn_abi, 0).unwrap_err(), - reference_param_validity(&source_types[1], &fn_abi, fn_abi.args.len()) + reference_param_validity(&instance, &source_types[1], &fn_abi, 0) .unwrap_err(), + reference_param_validity( + &instance, + &source_types[1], + &fn_abi, + fn_abi.args.len(), + ) + .unwrap_err(), ]; let mut wrong_mode = fn_abi.clone(); wrong_mode.args[1].mode = fn_abi.args[15].mode.clone(); assert!(matches!(wrong_mode.args[1].mode, PassMode::Indirect { .. })); errors.push( - reference_param_validity(&source_types[1], &wrong_mode, 1).unwrap_err(), + reference_param_validity(&instance, &source_types[1], &wrong_mode, 1) + .unwrap_err(), ); wrong_mode.args[3].mode = fn_abi.args[1].mode.clone(); errors.push( - reference_param_validity(&source_types[3], &wrong_mode, 3).unwrap_err(), + reference_param_validity(&instance, &source_types[3], &wrong_mode, 3) + .unwrap_err(), ); std::ops::ControlFlow::<(), _>::Continue((facts, errors)) }) diff --git a/cuda-oxide/crates/mir-importer/src/translator/types.rs b/cuda-oxide/crates/mir-importer/src/translator/types.rs index f64b01c999..e3ad890548 100644 --- a/cuda-oxide/crates/mir-importer/src/translator/types.rs +++ b/cuda-oxide/crates/mir-importer/src/translator/types.rs @@ -902,10 +902,8 @@ pub fn translate_type( // Total struct size (bytes) let size: u64 = shape.size.bytes() as u64; let abi_kind = if adt_def.repr().flags.is_transparent - && matches!( - &shape.value_repr, - rustc_public::abi::ValueRepr::Scalar(_) - ) { + && matches!(&shape.abi, rustc_public::abi::ValueAbi::Scalar(_)) + { StructAbiKind::TransparentScalar } else { StructAbiKind::Aggregate diff --git a/cuda-oxide/crates/rustc-codegen-cuda/README.md b/cuda-oxide/crates/rustc-codegen-cuda/README.md index 775c34399b..d4c9c405c6 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/README.md +++ b/cuda-oxide/crates/rustc-codegen-cuda/README.md @@ -229,6 +229,6 @@ cargo oxide pipeline - **Closures work**. Both `move` closures (capture by value) and non-move closures (capture by reference via HMM) can be passed to kernels. - **Unwind paths are unreachable**. The backend ignores all MIR unwind edges, so `panic=abort` and custom sysroots are unnecessary. If a panic condition is hit at runtime, the GPU traps. -### Nightly Rust +### Pinned Rust Toolchain -The backend uses `#![feature(rustc_private)]` and pins to a specific nightly via `rust-toolchain.toml`. The workspace currently uses `nightly-2026-10-11`. +The backend uses `#![feature(rustc_private)]` (unlocked on stable via `RUSTC_BOOTSTRAP=1`) and pins to a specific toolchain release via `rust-toolchain.toml`. The workspace currently uses `1.99.1`. diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/README.md b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/README.md index a99a7bdca3..eb9de3f9a1 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/README.md +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/README.md @@ -14,7 +14,7 @@ Atomic operations require **LLVM 22 or newer** for correct syncscope generation. Without it, scopes (`.gpu`, `.cta`, `.sys`) will be missing from generated PTX. -The pinned Rust toolchain (`nightly-2026-10-11`) ships LLVM 23 with NVPTX +The pinned Rust toolchain (`1.99.1`) ships LLVM 23 with NVPTX enabled via the `llvm-tools` component, so the default onboarding path already satisfies this requirement: @@ -83,7 +83,7 @@ CUDA_OXIDE_VERBOSE=1 cargo oxide run atomics ### Floating-point add lowering -With the pinned `nightly-2026-10-11` toolchain's LLVM 23 NVPTX defaults, +With the pinned `1.99.1` toolchain's LLVM 23 NVPTX defaults, `atomic_f32_fetch_add_test` uses a compare-and-swap (CAS) loop for global memory. Generic-address-space `f32` adds also use a loop; shared `f32` and `f64` adds still use native instructions on supported targets. For example: diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/src/main.rs index 774cd77142..ddaca5cd7a 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/src/main.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/atomics/src/main.rs @@ -545,19 +545,15 @@ mod kernels { pub fn core_atomic_ordering_probe(counter: *mut u32, mut out: DisjointSlice) { let gid = thread::index_1d(); let pointer = counter; - // nightly-2026-08-28 added a `const VOLATILE: bool` generic to the - // load/store intrinsics; plain atomics pass `false`. let current = unsafe { - core::intrinsics::atomic_load::( + core::intrinsics::atomic_load::( pointer, ) }; unsafe { - core::intrinsics::atomic_store::< - u32, - { core::intrinsics::AtomicOrdering::Release }, - false, - >(pointer, current) + core::intrinsics::atomic_store::( + pointer, current, + ) }; let swapped = unsafe { core::intrinsics::atomic_xchg::( diff --git a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs index 08e0784125..f6bee22ce3 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/examples/function_item_call/src/main.rs @@ -18,6 +18,8 @@ //! `call_once_decoy` proves that an ordinary function whose name contains //! trait-like text is not mistaken for a callable-trait shim. +#![feature(never_type)] + use cuda_core::simt::LaunchConfig; use cuda_core::{CudaContext, DeviceBuffer}; use cuda_device::{DisjointSlice, device, kernel, thread}; diff --git a/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs b/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs index 83ab7918bd..078099ebaf 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/src/collector.rs @@ -1524,7 +1524,7 @@ impl<'tcx> DeviceCollector<'tcx> { // kind (a reified `#[target_feature]` fn, say) must be checked. if !matches!( func.instance.def, - InstanceKind::Shim(ShimKind::DropGlue(..) | ShimKind::FnPtrAsPtr(..)) + InstanceKind::Shim(ShimKind::DropGlue(..) | ShimKind::FnPtrAddr(..)) ) { self.check_host_cpu_only(def_id, ctx.user_span, &ctx, None); } @@ -1540,7 +1540,7 @@ impl<'tcx> DeviceCollector<'tcx> { // `instance_mir` body: drop glue, and (since // nightly-2026-08-28) the `::as_ptr` // shim that backs `FnPtr::addr`. - InstanceKind::Shim(ShimKind::DropGlue(..) | ShimKind::FnPtrAsPtr(..)) + InstanceKind::Shim(ShimKind::DropGlue(..) | ShimKind::FnPtrAddr(..)) ); if has_mir { // Use instance_mir for monomorphized MIR. @@ -2054,7 +2054,7 @@ impl<'tcx> DeviceCollector<'tcx> { if !matches!( resolved.def, InstanceKind::Item(_) - | InstanceKind::Shim(ShimKind::DropGlue(..) | ShimKind::FnPtrAsPtr(..)) + | InstanceKind::Shim(ShimKind::DropGlue(..) | ShimKind::FnPtrAddr(..)) ) { return; } @@ -2178,7 +2178,7 @@ impl<'tcx> DeviceCollector<'tcx> { // `is_mir_available` is false for their def_id, but `instance_mir` // synthesises their body; do not skip those. if !self.tcx.is_mir_available(resolved.def_id()) - && !matches!(resolved.def, InstanceKind::Shim(ShimKind::FnPtrAsPtr(..))) + && !matches!(resolved.def, InstanceKind::Shim(ShimKind::FnPtrAddr(..))) { if self.verbose { eprintln!( @@ -2726,7 +2726,7 @@ impl<'tcx> DeviceCollector<'tcx> { .with_help( "use the cuda-intrinsics crate generated by this cuda-oxide revision, then rebuild the device crate", ) - .emit_fatal() + .emit() } /// Emits the heap-allocation diagnostic (issue #108) and aborts. @@ -2762,7 +2762,7 @@ impl<'tcx> DeviceCollector<'tcx> { "store the data in a fixed-size array `[T; N]` instead, or in a \ `SharedArray` when the scratch space should be shared by the thread block", ) - .emit_fatal() + .emit() } /// Refuses a device-reachable function that only makes sense on the host diff --git a/cuda-oxide/crates/rustc-codegen-cuda/src/device_codegen.rs b/cuda-oxide/crates/rustc-codegen-cuda/src/device_codegen.rs index f7e7ae0103..f22d0937d9 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/src/device_codegen.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/src/device_codegen.rs @@ -103,7 +103,6 @@ use rustc_middle::mir::interpret::{AllocId, GlobalAlloc, Scalar}; use rustc_middle::mir::visit::Visitor; use rustc_middle::mir::{ConstOperand, ConstValue, Location}; use rustc_middle::mono::MonoItem; -use rustc_middle::ty::consts::ConstExt; use rustc_middle::ty::layout::{LayoutCx, LayoutOf}; use rustc_middle::ty::{EarlyBinder, Instance, InstanceKind, ShimKind, TypingEnv}; use rustc_middle::ty::{Ty, TyCtxt, TyKind}; @@ -218,7 +217,7 @@ fn rustc_ty_to_device_extern_type<'tcx>( "failed to compute layout for repr(transparent) device-extern type `{ty}`: {err:?}" ) })?; - let peeled = layout.peel_transparent_wrappers_from_non_1zst(&layout_cx); + let peeled = layout.peel_transparent_wrappers(&layout_cx); // A transparent type with no peelable non-1ZST field is not a // scalar/pointer wrapper that this device-extern ABI can represent. diff --git a/cuda-oxide/crates/rustc-codegen-cuda/src/lib.rs b/cuda-oxide/crates/rustc-codegen-cuda/src/lib.rs index 4ea4e048fe..8564f84dbe 100644 --- a/cuda-oxide/crates/rustc-codegen-cuda/src/lib.rs +++ b/cuda-oxide/crates/rustc-codegen-cuda/src/lib.rs @@ -341,7 +341,7 @@ use rustc_middle::dep_graph::WorkProductMap; use rustc_middle::ty::TyCtxt; use rustc_middle::ty::print::with_no_trimmed_paths; use rustc_session::config::OutputFilenames; -use rustc_session::{CodegenBackendInit, EarlySession, IncrCompSession, Session}; +use rustc_session::{IncrCompSession, Session}; use std::any::Any; use std::collections::BTreeSet; use std::path::{Path, PathBuf}; @@ -562,12 +562,12 @@ impl CodegenBackend for CudaCodegenBackend { "cuda" } - fn init(&mut self, sess: &EarlySession) -> CodegenBackendInit { + fn init(&self, sess: &Session) { // Note: Don't log here - init() is called for ALL crates including dependencies. // We log in codegen_crate() only when there are kernels to compile. // Initialize the underlying LLVM backend - self.llvm_backend.init(sess) + self.llvm_backend.init(sess); } fn print_version(&self) { @@ -582,7 +582,7 @@ impl CodegenBackend for CudaCodegenBackend { self.llvm_backend.target_cpu(sess) } - fn target_config(&self, sess: &EarlySession) -> rustc_codegen_ssa::TargetConfig { + fn target_config(&self, sess: &Session) -> rustc_codegen_ssa::TargetConfig { self.llvm_backend.target_config(sess) } @@ -966,6 +966,7 @@ impl CodegenBackend for CudaCodegenBackend { assembly: None, llvm_ir: None, global_asm_object: None, + links_from_incr_cache: Vec::new(), }); } (compiled_modules, work_products) diff --git a/cuda-oxide/cuda-oxide-book/appendix/building-from-source.md b/cuda-oxide/cuda-oxide-book/appendix/building-from-source.md index 56500c646a..13f9b29936 100644 --- a/cuda-oxide/cuda-oxide-book/appendix/building-from-source.md +++ b/cuda-oxide/cuda-oxide-book/appendix/building-from-source.md @@ -10,7 +10,7 @@ from a fresh checkout. If you just want to run an example, the | Dependency | Version | Purpose | |:-----------------|:----------------------------- |:------------------------------------------------------------| -| **Rust nightly** | `nightly-2026-10-11` (pinned) | Compiler toolchain with `rustc-dev` for the codegen backend | +| **Rust** | `1.99.1` (pinned stable) | Compiler toolchain with `rustc-dev` for the codegen backend | | **CUDA Toolkit** | 13.0+ (with cuRAND headers) | Driver API, `nvcc`, PTX assembler; `curand.h` for bindgen | | **Clang** | 21+ (`clang-21` pkg) | `bindgen` in host `cuda-bindings` needs clang's headers | | **Linux** | Tested on Ubuntu 24.04 | Windows and macOS are not supported | @@ -26,22 +26,21 @@ cd cuda-rust ## Install the Rust toolchain The CUDA Oxide component ships `cuda-oxide/rust-toolchain.toml`, which pins the -exact nightly version and components. Rustup picks it up automatically for -commands run inside `cuda-oxide/`; the repository root and `cutile-rs/` stay on -stable: +exact stable release and components. Rustup picks it up automatically for +commands run inside `cuda-oxide/`: ```toml # cuda-oxide/rust-toolchain.toml [toolchain] -channel = "nightly-2026-10-11" +channel = "1.99.1" components = ["rust-src", "rustc-dev", "rust-analyzer", "clippy", "rustfmt", "llvm-tools"] ``` If you need to install manually: ```bash -rustup toolchain install nightly-2026-10-11 -rustup component add rust-src rustc-dev rust-analyzer clippy rustfmt llvm-tools --toolchain nightly-2026-10-11 +rustup toolchain install 1.99.1 +rustup component add rust-src rustc-dev rust-analyzer clippy rustfmt llvm-tools --toolchain 1.99.1 ``` `rust-src` provides the standard library source for cross-compilation, @@ -66,7 +65,7 @@ required for `ptxas` and header files, but you will not be able to run kernels. ## Install LLVM (usually optional) The codegen pipeline emits LLVM IR and invokes `llc` to produce PTX. The -pinned Rust toolchain (`nightly-2026-10-11`) already ships LLVM 23 with the +pinned Rust toolchain (`1.99.1`) already ships LLVM 23 with the NVPTX backend enabled via the `llvm-tools` component, so the recommended path is: @@ -156,10 +155,10 @@ build process. `cargo-oxide` handles building it transparently. ## Install cargo-oxide `cargo-oxide` is the cargo subcommand that drives the full compilation -pipeline. Inside the repo, it works via a workspace alias. For standalone use, install it with the pinned nightly toolchain: +pipeline. Inside the repo, it works via a workspace alias. For standalone use, install it with the pinned toolchain: ```bash -cargo +nightly-2026-10-11 install --git https://github.com/NVIDIA/cuda-rust.git cargo-oxide +cargo +1.99.1 install --git https://github.com/NVIDIA/cuda-rust.git cargo-oxide ``` On first run, `cargo-oxide` automatically fetches and builds the codegen backend @@ -287,7 +286,7 @@ repository/ `error[E0463]: can't find crate for rustc_middle` : You are missing the `rustc-dev` component. Run: - `rustup component add rustc-dev --toolchain nightly-2026-10-11`. + `rustup component add rustc-dev --toolchain 1.99.1`. CUDA driver version mismatch : The toolkit version (compile-time) and driver version (runtime) must be diff --git a/cuda-oxide/cuda-oxide-book/compiler/images/dev-workflow-shift.svg b/cuda-oxide/cuda-oxide-book/compiler/images/dev-workflow-shift.svg new file mode 100644 index 0000000000..eefdafc85e --- /dev/null +++ b/cuda-oxide/cuda-oxide-book/compiler/images/dev-workflow-shift.svg @@ -0,0 +1,107 @@ + + + +Track 1: Shifting cuda-oxide development from a daily nightly pin to an official stable Rust release + + + + + + + + +Track 1: Moving Development & Repo Toolchains to Stable Rust +Pins an official stable Rust release (`1.99.1`) in `cuda-oxide/rust-toolchain.toml` while keeping `rustc_private` and `rustc_public`. + + +BEFORE · DAILY NIGHTLY PIN IN CUDA-OXIDE + +Nightly Pin in cuda-oxide/rust-toolchain.toml + +Root + cutile-rs: channel = "1.98.0" (Stable) +cuda-oxide/: channel = "nightly-2026-10-11" +Result: daily snapshot drift and mid-cycle rustc churn. + + + + +Why Nightly Was Used Originally + +• Needed #![feature(rustc_private)] for rustc-codegen-cuda +• Needed rustc_public (stable_mir) for mir-importer +• Needed -Z codegen-backend and -Z mir-enable-passes + + +NEW WORKFLOW · PINNED STABLE RELEASE (1.99.1) + + +Step 1: Pin Stable Toolchain + + +channel = "1.99.1" # Stable! +components = ["rustc-dev", ...] +Ships LLVM 23.1.1 with NVPTX. +librustc_public.rlib included! + + + + +Step 2: cargo-oxide Sets Env + + +cmd.env("RUSTC_BOOTSTRAP", "1") +-Z codegen-backend=... +Works for both building the .so +and compiling kernel crates. + + + + +Step 3: 6-Week Release Cadence + + +1.99.1 -> 1.100.0 -> 1.101.0 +Vetted through 6 weeks of beta +No daily nightly regressions. +Same model as Microsoft UTC. + diff --git a/cuda-oxide/cuda-oxide-book/compiler/images/ecosystem-and-bootstrap.svg b/cuda-oxide/cuda-oxide-book/compiler/images/ecosystem-and-bootstrap.svg new file mode 100644 index 0000000000..f629332803 --- /dev/null +++ b/cuda-oxide/cuda-oxide-book/compiler/images/ecosystem-and-bootstrap.svg @@ -0,0 +1,110 @@ + + + +How custom Rust backends ship in production and why stable rustc already contains everything cuda-oxide needs + + + + + + + + +The Ecosystem Playbook & Why Stable rustc Already Has Everything We Need +Official stable Rust releases ship the exact same backend hooks and rustc_public crates as nightly behind RUSTC_BOOTSTRAP=1. + + +1 · WHAT OTHER RUST COMPILER / BACKEND PROJECTS DO + + +A. Reimplement Frontend (rust-analyzer, cutile-rs) +• Zero rustc code; custom parser + type solver on stable. +• Works for IDEs, not codegen (no borrowck, mono, or LLVM). + + +B. Out-of-Tree Nightly Pin (rust-gpu, Kani, Charon) +• Uses #![feature(rustc_private)] on a daily nightly pin. +• Pain point: users need nightly; daily builds can hit ICEs. + + +C. In-Tree Upstream (rustc_codegen_cranelift / gcc) +• Lives inside rust-lang/rust; built by upstream Rust CI. +• Only possible if upstream Rust owns your target backend. + + +D. Stable-Tag SDK Bundle (MS codegen_utc, Ferrocene) +• Uses the exact same CodegenBackend trait we use! +• Builds against a vetted STABLE release (e.g. 1.90 / 1.99) + and ships rustc + .so + sysroot as one standalone bundle. + + +2 · WHAT IS ACTUALLY INSIDE AN OFFICIAL STABLE RUST RELEASE + + +Official rustup Stable Toolchain (e.g. 1.99.1) + + + +Startup Gatekeeper in rustc +if channel == "nightly" || RUSTC_BOOTSTRAP == "1" + + +Unlocks on Stable: + + + +1. Unstable Compiler Flags (-Z) +-Z codegen-backend=... -Z mir-enable-passes=... + + +2. rustc-dev Component (Ships on Stable rustup!) +librustc_driver.so + librustc_public.rlib +(Verified: rustc_public & bridge come for free!) + + +3. Scoped Feature Gate (-Z allow-features=...) +Unlocks rustc_private for us; user code stays stable. + diff --git a/cuda-oxide/cuda-oxide-book/compiler/images/toolkit-bundle-and-sizes.svg b/cuda-oxide/cuda-oxide-book/compiler/images/toolkit-bundle-and-sizes.svg new file mode 100644 index 0000000000..f0e96380ab --- /dev/null +++ b/cuda-oxide/cuda-oxide-book/compiler/images/toolkit-bundle-and-sizes.svg @@ -0,0 +1,113 @@ + + + +Track 2: Self-contained CUDA Toolkit bundle, measured binary sizes, and bug-fix workflow + + + + + + + + +Track 2: Bundling Everything for Users & the CUDA Toolkit (Option B) +Users never install rustc-dev or compile the backend. Everything ships prebuilt in one ~120 MB compressed package. + + +1 · MEASURED SIZE BREAKDOWN (x86_64-unknown-linux-gnu) + + +Stays in NVIDIA CI Only (Not Shipped): ~1.15 GB + +• rustc-dev (librustc_*.rlib compiler libs): ~930 MB +• llvm-tools & unstripped debug artifacts: ~220 MB + + +CI builds & strips bundle + + +Shipped in CUDA Toolkit / Standalone Bundle + + +librustc_codegen_cuda.so +34 MB (8.6 MB zst) + + +librustc_driver.so (frontend) +159 MB (~42 MB zst) + + +libLLVM.so (host LLVM backend) +199 MB (~52 MB zst) + + +Pre-encoded sysroot + bins +61 MB (~17 MB zst) + +Total Bundle: ~453 MB on disk | ~120 MB compressed + + +2 · ZERO-CONFIG USER FLOW & BUG OWNERSHIP + + +User Experience: Zero rustup or Nightly Setup + +$ cargo oxide build +1. cargo-oxide points RUSTC at bundled cuda-oxide-rustc +2. Uses bundled sysroot (built with -Z always-encode-mir) +3. Emits host binary + embedded PTX/CUBIN in one pass + + +How We Fix Bugs Without Maintaining a rustc Fork + + +Bucket 1 (~95% of bugs): In cuda-oxide crates +Fix in mir-importer/lower -> ship 8.6 MB .so hotfix only. + + +Bucket 2 (~5% of bugs): Upstream rustc passes +Toggle pass flag in driver (-JumpThreading) or bump tag. + diff --git a/cuda-oxide/cuda-oxide-book/compiler/images/upstream-rustc-public-flow.svg b/cuda-oxide/cuda-oxide-book/compiler/images/upstream-rustc-public-flow.svg new file mode 100644 index 0000000000..c0eecfc213 --- /dev/null +++ b/cuda-oxide/cuda-oxide-book/compiler/images/upstream-rustc-public-flow.svg @@ -0,0 +1,104 @@ + + + +Upstreaming rustc_public PRs on a stable release cycle without waiting or emergency nightly bumps + + + + + + + + +Upstreaming to rustc_public Without Waiting for the 6-Week Stable Train +Because rustc_public wraps existing rustc_middle data, facts.rs lets us ship on Day 1 and switch to rustc_public on schedule. + + +1 · OLD NIGHTLY WORKFLOW (EXAMPLE: PR #1482 + PR #1483) + + +1. Upstream PR to Rust + +rust-lang/rust#163242 +Exposes ArgAttributes getters + + + + +2. Forced Nightly Bump (PR #1482) + +nightly-08-28 -> nightly-10-11 +1,506 files changed due to unrelated churn! + + + + +3. Land Feature (PR #1483) + +facts.rs calls attrs.regular() +Blocked until step 2 finished + + +2 · NEW STABLE WORKFLOW (ZERO WAIT, ZERO EMERGENCY BUMPS) + + +Day 1: Ship Feature Immediately on Current Stable + + +// Inside facts.rs (quarantine boundary): +let internal = rustc_internal::internal(tcx, arg); +• Underlying rustc_abi data ALREADY exists in current stable! +• Upstream rust#163242 in parallel; 0 toolchain churn. + + +6 wks + + +Week 6: Scheduled Stable Release Upgrade + + +// Swap 5 lines in facts.rs to rustc_public: +let flags = attrs.regular(); +• rust#163242 arrives automatically in the next stable tag. +• Delete the 5-line shim in facts.rs; rest is unchanged. + diff --git a/cuda-oxide/cuda-oxide-book/compiler/lowering-pipeline.md b/cuda-oxide/cuda-oxide-book/compiler/lowering-pipeline.md index e227a9e59c..668013f72a 100644 --- a/cuda-oxide/cuda-oxide-book/compiler/lowering-pipeline.md +++ b/cuda-oxide/cuda-oxide-book/compiler/lowering-pipeline.md @@ -587,7 +587,7 @@ that cuda-oxide emits. | 5th | `llc-21` on `PATH` | Distro / `apt.llvm.org` install of LLVM 21. | | 6th | `llc` on `PATH` | Reporting fallback only; rejected at runtime if older than LLVM 21. | -The pinned Rust toolchain (`nightly-2026-10-11`) ships LLVM 23 with NVPTX +The pinned Rust toolchain (`1.99.1`) ships LLVM 23 with NVPTX enabled, so `rustup component add llvm-tools` is the recommended onboarding path. The PATH probes for `llc-23` / `llc-22` / `llc-21` are kept as a fallback for users with an existing LLVM install. If none of the probes succeed the diff --git a/cuda-oxide/cuda-oxide-book/compiler/rustc-public.md b/cuda-oxide/cuda-oxide-book/compiler/rustc-public.md index b157eda3d1..d046635d95 100644 --- a/cuda-oxide/cuda-oxide-book/compiler/rustc-public.md +++ b/cuda-oxide/cuda-oxide-book/compiler/rustc-public.md @@ -277,42 +277,91 @@ already explicit. The `mir-importer` crate translates this into Pliron IR (an MLIR-like framework), and from there the [lowering pipeline](lowering-pipeline.md) takes it the rest of the way to PTX. -## Nightly pinning +## Stabilizing the toolchain pin -Even with `rustc_public` providing a stable *API*, the bridge layer -(`rustc_public_bridge`) is still compiled against the compiler's internals -and is not independently versioned. In practice, this means a specific -version of cuda-oxide works with a specific nightly. - -cuda-oxide pins to an exact nightly release via `rust-toolchain.toml`: +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 +`rust-toolchain.toml`: ```toml [toolchain] -channel = "nightly-2026-10-11" +channel = "1.99.1" components = ["rust-src", "rustc-dev", "rust-analyzer", "clippy", "rustfmt", "llvm-tools"] ``` -This pin guarantees reproducible builds: anyone who clones the repository -gets the same compiler, the same MIR shapes, and the same bridge behavior. -When updating the pin, the process is: - -1. Bump the nightly date in `cuda-oxide/rust-toolchain.toml`. Rustup walks up - from the standalone `rustc-codegen-cuda` workspace and finds that component - pin, so no nested copy is needed. `cuda-oxide/scripts/check-toolchain-parity.sh` - checks the scaffold `cargo oxide new` writes, the devcontainer, and the - blocks and dated commands quoted in this book and the READMEs -- run it to - find the stragglers. -2. Fix any `rustc_public` API changes (usually minor -- that is the whole - point of the stable API). -3. Run the full test suite to verify that all examples still compile and - produce correct PTX. -4. Celebrate, or revert and try next week's nightly. +```{figure} images/ecosystem-and-bootstrap.svg +:align: center +:width: 100% -```{note} -As `rustc_public` matures and moves toward a crates.io release, the coupling -to a specific nightly will loosen. The long-term goal is for cuda-oxide to -work with any sufficiently recent stable Rust toolchain -- but we are not -there yet. +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` + +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 +`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`, +`-Z always-encode-mir`) on the pinned stable compiler while giving developers +and users a compiler that has already completed Rust's 12-week nightly and beta +stabilization cycle. + +```{figure} images/dev-workflow-shift.svg +: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`. +``` + +### 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 +`rust-lang/rust#163242`), it takes 6 to 12 weeks for that commit to reach the +next stable Rust release. + +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: + +1. **Day 1 (on the current stable pin):** Land the feature immediately by + reading the underlying `rustc_middle` / `rustc_target` field inside + `facts.rs` via `rustc_internal::internal(tcx, ...)`, while upstreaming the + `rustc_public` getter in parallel. +2. **Next scheduled stable upgrade:** Once the new stable Rust release ships + with the upstreamed `rustc_public` API, replace the temporary internal read + in `facts.rs` with the `rustc_public` method call. + +```{figure} images/upstream-rustc-public-flow.svg +:align: center +:width: 100% + +Decoupling feature delivery from toolchain bumps using `facts.rs` and scheduled stable upgrades. +``` + +### Self-contained bundling for standalone and CUDA Toolkit distribution + +For standalone or CUDA Toolkit packaging (following the same model as +Microsoft's `rustc_codegen_utc` and Ferrocene), users do not need `rustc-dev` +(~930 MB of compiler-internal `.rlib`s used only when building the backend +itself). A stripped runtime bundle containing `rustc`, `cargo`, `cargo-oxide`, +`librustc_codegen_cuda.so` (34 MB stripped, 8.6 MB compressed), +`librustc_driver.so`, `libLLVM.so`, and a pre-encoded `core`/`alloc`/`std` +sysroot totals roughly 453 MB on disk (~120 MB compressed). + +```{figure} images/toolkit-bundle-and-sizes.svg +:align: center +:width: 100% + +Measured bundle footprint on `x86_64-unknown-linux-gnu` and the two-bucket maintenance workflow. ``` --- diff --git a/cuda-oxide/cuda-oxide-book/getting-started/hello-gpu.md b/cuda-oxide/cuda-oxide-book/getting-started/hello-gpu.md index 3a4ff4ec5e..61793a456e 100644 --- a/cuda-oxide/cuda-oxide-book/getting-started/hello-gpu.md +++ b/cuda-oxide/cuda-oxide-book/getting-started/hello-gpu.md @@ -6,10 +6,10 @@ This section walks through installing cuda-oxide, creating a project, writing a ## Install cargo-oxide -If you haven't already, install the build tool with the pinned nightly toolchain: +If you haven't already, install the build tool with the pinned toolchain: ```bash -cargo +nightly-2026-10-11 install --git https://github.com/NVIDIA/cuda-rust.git cargo-oxide +cargo +1.99.1 install --git https://github.com/NVIDIA/cuda-rust.git cargo-oxide ``` Verify that your environment is set up correctly: diff --git a/cuda-oxide/cuda-oxide-book/getting-started/installation.md b/cuda-oxide/cuda-oxide-book/getting-started/installation.md index d495e557c0..f8114dcc48 100644 --- a/cuda-oxide/cuda-oxide-book/getting-started/installation.md +++ b/cuda-oxide/cuda-oxide-book/getting-started/installation.md @@ -279,15 +279,15 @@ Validated on Asus GX10 / NVIDIA DGX Spark. ## Rust toolchain The CUDA Oxide workspace ships `cuda-oxide/rust-toolchain.toml`, which pins the -exact nightly version and required components. When you first run a `cargo` +exact stable release and required components. When you first run a `cargo` command inside `cuda-oxide/`, rustup installs the correct toolchain -automatically. Commands at the repository root and in `cutile-rs/` use stable. +automatically. If you need to install it manually: ```bash -rustup toolchain install nightly-2026-10-11 -rustup component add rust-src rustc-dev rust-analyzer clippy rustfmt llvm-tools --toolchain nightly-2026-10-11 +rustup toolchain install 1.99.1 +rustup component add rust-src rustc-dev rust-analyzer clippy rustfmt llvm-tools --toolchain 1.99.1 ``` Three of these components are what the codegen backend and doctor need: @@ -306,10 +306,10 @@ The other three are not needed to build: `rust-analyzer` powers IDE support, `cl **Inside the cuda-oxide repo**, it works out of the box via a workspace alias -- no extra install step. -**For use outside the repo** (your own projects), install it with the pinned nightly toolchain: +**For use outside the repo** (your own projects), install it with the pinned toolchain: ```bash -cargo +nightly-2026-10-11 install --git https://github.com/NVIDIA/cuda-rust.git cargo-oxide +cargo +1.99.1 install --git https://github.com/NVIDIA/cuda-rust.git cargo-oxide ``` On first run, `cargo-oxide` will automatically fetch and build the codegen backend. Subsequent runs reuse the cached build. @@ -430,8 +430,8 @@ If everything is configured correctly, this compiles a Rust kernel to PTX, launc :::{tip} **Common issues:** -- `No working llc found` -- prefer `rustup component add llvm-tools --toolchain nightly-2026-10-11`, or install LLVM 21+ (`sudo apt install llvm-21`), add `/usr/lib/llvm-21/bin` to your `PATH`, or set `CUDA_OXIDE_LLC=/usr/bin/llc-21`. +- `No working llc found` -- prefer `rustup component add llvm-tools --toolchain 1.99.1`, or install LLVM 21+ (`sudo apt install llvm-21`), add `/usr/lib/llvm-21/bin` to your `PATH`, or set `CUDA_OXIDE_LLC=/usr/bin/llc-21`. - `'stddef.h' file not found` when building host `cuda-bindings` -- install clang dev headers: `sudo apt install clang-21` (or `libclang-common-21-dev`). - `cuda.h not found` -- Set `CUDA_TOOLKIT_PATH` to your CUDA install root, or ensure `/usr/local/cuda/include/cuda.h` exists. -- `rust-src` / `llvm-tools` component missing -- Run `rustup component add rust-src llvm-tools --toolchain nightly-2026-10-11`. +- `rust-src` / `llvm-tools` component missing -- Run `rustup component add rust-src llvm-tools --toolchain 1.99.1`. ::: diff --git a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_accumulator_tile_rows.stderr b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_accumulator_tile_rows.stderr index 131426e346..1b2cf9f02a 100644 --- a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_accumulator_tile_rows.stderr +++ b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_accumulator_tile_rows.stderr @@ -2,6 +2,6 @@ error[E0599]: no method named `mma_mxf4_pair` found for struct `AccumulatorTile1 --> tests/compile_fail/primitives_accumulator_tile_rows.rs:11:14 | 11 | tile.mma_mxf4_pair::<0>( - | ^^^^^^^^^^^^^ method not found in `AccumulatorTile16x8<1, 8>` + | -----^^^^^^^^^^^^^ method not found in `AccumulatorTile16x8<1, 8>` | = note: the method was found for `AccumulatorTile16x8<2, COLS>` diff --git a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_bulk_wait_count.stderr b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_bulk_wait_count.stderr index c294f40ce4..54bef41b69 100644 --- a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_bulk_wait_count.stderr +++ b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_bulk_wait_count.stderr @@ -17,6 +17,7 @@ note: required by a bound in `cp_async_bulk_wait_group_read` --> src/tma.rs | | pub fn cp_async_bulk_wait_group_read() + | ----------------------------- required by a bound in this function | where | (): ValidWaitCount, | ^^^^^^^^^^^^^^^^^ required by this bound in `cp_async_bulk_wait_group_read` diff --git a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_mma_destination.stderr b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_mma_destination.stderr index 6e600320ac..0e624927fe 100644 --- a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_mma_destination.stderr +++ b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_mma_destination.stderr @@ -12,5 +12,6 @@ note: required by a bound in `ctm_rs::tcgen05::mma::mma` --> src/tcgen05/mma.rs | | pub unsafe fn mma( + | --- required by a bound in this function | dst: &impl TmemOperand, | ^^^^^^^^^^^^^^^^^^^^^ required by this bound in `mma` diff --git a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_mma_k_step.stderr b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_mma_k_step.stderr index 9312efd2dc..11de33be85 100644 --- a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_mma_k_step.stderr +++ b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_mma_k_step.stderr @@ -19,9 +19,8 @@ help: `()` implements trait `ValidKStep` note: required by a bound in `ctm_rs::tcgen05::descriptors::SmemOperand::<'s, R, F>::k16` --> src/tcgen05/descriptors.rs | - | impl<'s, R: OperandRole, F: MmaInputFormat> SmemOperand<'s, R, F> { -... | pub fn k16(&self) -> Self + | --- required by a bound in this associated function | where | (): ValidKStep, | ^^^^^^^^^^^^^^^^ required by this bound in `SmemOperand::<'s, R, F>::k16` diff --git a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_register_count_alignment.stderr b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_register_count_alignment.stderr index e599980ea1..0f732206a4 100644 --- a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_register_count_alignment.stderr +++ b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_register_count_alignment.stderr @@ -39,6 +39,7 @@ note: required by a bound in `setmaxnreg_dec` --> src/lib.rs | | pub unsafe fn setmaxnreg_dec() + | -------------- required by a bound in this function | where | (): ValidRegisterCount, | ^^^^^^^^^^^^^^^^^^^^^ required by this bound in `setmaxnreg_dec` diff --git a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_register_count_range.stderr b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_register_count_range.stderr index b4099d8753..022ee950f2 100644 --- a/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_register_count_range.stderr +++ b/cuda-oxide/cute-rs-preview/ctm-rs/tests/compile_fail/primitives_register_count_range.stderr @@ -39,6 +39,7 @@ note: required by a bound in `setmaxnreg_inc` --> src/lib.rs | | pub unsafe fn setmaxnreg_inc() + | -------------- required by a bound in this function | where | (): ValidRegisterCount, | ^^^^^^^^^^^^^^^^^^^^^ required by this bound in `setmaxnreg_inc` diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/bf16_invalid_storage.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/bf16_invalid_storage.stderr index 0989d02e0f..cf8452f452 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/bf16_invalid_storage.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/bf16_invalid_storage.stderr @@ -15,6 +15,7 @@ note: required by a bound in `Tensor` --> src/tensor/mod.rs | | pub struct Tensor<'a, Element, Layout, Storage = Element> + | ------ required by a bound in this struct | where | Element: TensorElement, | ^^^^^^^^^^^^^^^^^^^^^^ required by this bound in `Tensor` diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/global_store_storage.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/global_store_storage.stderr index e445476fc2..21ec26f822 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/global_store_storage.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/global_store_storage.stderr @@ -20,3 +20,24 @@ note: required by a bound in `GlobalStoreTile` | | pub struct GlobalStoreTile( | ^^^^^^^^^^^^^^^^^^ required by this bound in `GlobalStoreTile` + +error[E0277]: the trait bound `f64: GlobalStoreElement` is not satisfied + --> tests/compile_fail/global_store_storage.rs:7:13 + | +7 | let _ = cute_rs::GlobalStoreTile([0.0f64; 4]); + | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ the trait `GlobalStoreElement` is not implemented for `f64` + | +help: the following other types implement trait `GlobalStoreElement` + --> src/copy/vector.rs + | + | impl GlobalStoreElement for f16 {} + | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ `f16` + | impl GlobalStoreElement for f32 {} + | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ `f32` + | impl GlobalStoreElement for u32 {} + | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ `u32` +note: required by a bound in `GlobalStoreTile` + --> src/copy/vector.rs + | + | pub struct GlobalStoreTile( + | ^^^^^^^^^^^^^^^^^^ required by this bound in `GlobalStoreTile` diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/invalid_storage.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/invalid_storage.stderr index 9b925b827d..0c20794013 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/invalid_storage.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/invalid_storage.stderr @@ -15,6 +15,7 @@ note: required by a bound in `Tensor` --> src/tensor/mod.rs | | pub struct Tensor<'a, Element, Layout, Storage = Element> + | ------ required by a bound in this struct | where | Element: TensorElement, | ^^^^^^^^^^^^^^^^^^^^^^ required by this bound in `Tensor` diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/packed_decode_callback_escape.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/packed_decode_callback_escape.stderr index 1984d481c3..2fb1d17433 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/packed_decode_callback_escape.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/packed_decode_callback_escape.stderr @@ -4,6 +4,6 @@ error[E0521]: borrowed data escapes outside of closure 8 | let mut escaped = None; | ----------- `escaped` declared here, outside of the closure body 9 | PackedDecode::with_validated(&storage, &storage, 1, 16, 8, 16, 16, |proof| { - | ----- `proof` is only valid in the closure body + | ----- `proof` is a reference that is only valid in the closure body 10 | escaped = Some(proof); | ^^^^^^^^^^^^^^^^^^^^^ `proof` escapes the closure body here diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_invalid_layout.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_invalid_layout.stderr index 5323792824..5e352eb438 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_invalid_layout.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_invalid_layout.stderr @@ -1,10 +1,6 @@ error[E0080]: evaluation panicked: assertion failed: contiguous >= 32 && contiguous.is_multiple_of(32) --> src/attention/sm100/shared.rs | - | impl SharedLayout { - | /// Reject unsupported contiguous swizzle spans and outer tile extents. - | const VALID: () = { -... | assert!(contiguous >= 32 && contiguous.is_multiple_of(32)); | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ evaluation of `cute_rs::attention::sm100::SharedLayout::<128, 16>::VALID` failed here diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_tile_array_extent.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_tile_array_extent.stderr index 225b540376..e1bf6b751e 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_tile_array_extent.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_tile_array_extent.stderr @@ -1,9 +1,6 @@ error[E0080]: evaluation panicked: assertion failed: TILES > 0 --> src/attention/sm100/shared.rs | - | impl SharedTileArray { -... - | const VALID: () = { | assert!(TILES > 0); | ^^^^^^^^^^^^^^^^^^ evaluation of `cute_rs::attention::sm100::SharedTileArray::, cute_rs::OperandB, 0>::VALID` failed here @@ -16,9 +13,5 @@ note: erroneous constant encountered error[E0080]: evaluation panicked: assertion failed: TILES <= u32::MAX / Self::TILE_BYTES --> src/attention/sm100/shared.rs | - | impl SharedTileArray { -... - | const VALID: () = { - | assert!(TILES > 0); | assert!(TILES <= u32::MAX / Self::TILE_BYTES); | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ evaluation of `cute_rs::attention::sm100::SharedTileArray::, cute_rs::OperandB, u32::MAX>::VALID` failed here diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_ts_k16.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_ts_k16.stderr index 42d64cc7b5..f283248157 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_ts_k16.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_attention_ts_k16.stderr @@ -1,10 +1,6 @@ error[E0080]: evaluation panicked: assertion failed: K == 32 || (K > 0 && K.is_multiple_of(64)) --> src/attention/sm100/mma.rs | - | impl MmaTs { - | /// Check the group-one TMEM/shared profile, supported K span and B layout shape. - | const VALID: () = { -... | assert!(K == 32 || (K > 0 && K.is_multiple_of(64))); | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ evaluation of `cute_rs::attention::sm100::MmaTs::, 128, 128, 16, 1>::VALID` failed here diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_mma_shape_mismatch.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_mma_shape_mismatch.stderr index c6b0d4426d..73a9c66de9 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_mma_shape_mismatch.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/sm100_mma_shape_mismatch.stderr @@ -1,10 +1,6 @@ error[E0080]: evaluation panicked: assertion failed: B0L::SHAPE.0 == (N0 / 2) as i64 && B0L::SHAPE.1 == K as i64 --> src/sm100/mma.rs | - | impl< -... - | const VALID: () = { -... | assert!(B0L::SHAPE.0 == (N0 / 2) as i64 && B0L::SHAPE.1 == K as i64); | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ evaluation of `cute_rs::Sm100TiledMma::, 0, cute_rs::RowMajor<128, 64>>, cute_rs::ComposedLayout, 0, cute_rs::RowMajor<80, 64>>, cute_rs::ComposedLayout, 0, cute_rs::RowMajor<96, 64>>, 192, 160>::VALID` failed here diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_composed_overflow.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_composed_overflow.stderr index 1d68fe6674..15d2a10c90 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_composed_overflow.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_composed_overflow.stderr @@ -1,9 +1,6 @@ error[E0080]: evaluation panicked: static layout offset exceeds i64 --> src/static_layout.rs | - | impl StaticLayout -... - | const MAX_OFFSET: usize = | swizzled_bound(checked_offset_add(Outer::MAX_OFFSET, parameter(O)), B, M, S); | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ evaluation of `, 1, cute_rs::Layout, cute_rs::C>> as cute_rs::static_layout::StaticLayout>::MAX_OFFSET` failed inside this call | diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_leaf_overflow.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_leaf_overflow.stderr index 3061e916cf..2d33e2e32e 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_leaf_overflow.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_leaf_overflow.stderr @@ -1,10 +1,6 @@ error[E0080]: evaluation panicked: static layout physical offset overflow --> src/static_layout.rs | - | impl StaticLayout for Layout, C> { -... - | const MAX_OFFSET: usize = { - | assert!(N > 0 && S >= 0); | checked_offset_mul(parameter(N - 1), parameter(S)) | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ evaluation of `, cute_rs::C> as cute_rs::static_layout::StaticLayout>::MAX_OFFSET` failed inside this call | diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_sum_overflow.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_sum_overflow.stderr index 32c0d534c9..7912c1777e 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_sum_overflow.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/static_layout_sum_overflow.stderr @@ -1,9 +1,6 @@ error[E0080]: evaluation panicked: static layout offset exceeds i64 --> src/static_layout.rs | - | impl StaticLayout for Layout<(A, B), (SA, SB)> -... - | const MAX_OFFSET: usize = | checked_offset_add(Layout::::MAX_OFFSET, Layout::::MAX_OFFSET); | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ evaluation of `, cute_rs::C<2>), (cute_rs::C, cute_rs::C<1>)> as cute_rs::static_layout::StaticLayout>::MAX_OFFSET` failed inside this call | diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/strided_tensor_storage_mismatch.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/strided_tensor_storage_mismatch.stderr index 63f3379748..07bc48eb01 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/strided_tensor_storage_mismatch.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/strided_tensor_storage_mismatch.stderr @@ -15,6 +15,29 @@ note: required by a bound in `StridedTensor` --> src/tensor/partition.rs | | pub struct StridedTensor<'a, Element, Layout = Strided1D, Storage = Element> + | ------------- required by a bound in this struct + | where + | Element: TensorElement, + | ^^^^^^^^^^^^^^^^^^^^^^ required by this bound in `StridedTensor` + +error[E0277]: the trait bound `BFloat16: TensorElement` is not satisfied + --> tests/compile_fail/strided_tensor_storage_mismatch.rs:11:73 + | +11 | let _invalid: Option> = None; + | ^^^^ the trait `TensorElement` is not implemented for `BFloat16` + | +help: the trait `TensorElement` is not implemented for `BFloat16` + but trait `TensorElement` is implemented for it + --> src/numeric.rs + | + | impl TensorElement for BFloat16 {} + | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + = help: for that trait implementation, expected `BFloat16`, found `f16` +note: required by a bound in `StridedTensor` + --> src/tensor/partition.rs + | + | pub struct StridedTensor<'a, Element, Layout = Strided1D, Storage = Element> + | ------------- required by a bound in this struct | where | Element: TensorElement, | ^^^^^^^^^^^^^^^^^^^^^^ required by this bound in `StridedTensor` diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_cursor_zero_stages.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_cursor_zero_stages.stderr index 835a810591..c6ad8e625e 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_cursor_zero_stages.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_cursor_zero_stages.stderr @@ -1,8 +1,6 @@ error[E0080]: evaluation panicked: assertion failed: STAGES > 0 --> src/attention/sm100/pipeline.rs | - | impl RingCursor { - | /// Reject zero-stage cursors before any slot arithmetic. | const VALID: () = assert!(STAGES > 0); | ^^^^^^^^^^^^^^^^^^^ evaluation of `cute_rs::attention::sm100::pipeline::RingCursor::<0, cute_rs::Producer>::VALID` failed here diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_non_tc_wrong_signal.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_non_tc_wrong_signal.stderr index 13d70336d0..831db8f291 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_non_tc_wrong_signal.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_non_tc_wrong_signal.stderr @@ -15,9 +15,8 @@ help: the following other types implement trait `SignalPolicy` note: required by a bound in `TwoCtaRing::::thread_empty` --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { -... | pub unsafe fn thread_empty(self, ticket: u32) + | ------------ required by a bound in this associated function | where | Empty: SignalPolicy, | ^^^^^^^^^^^^ required by this bound in `TwoCtaRing::::thread_empty` @@ -36,9 +35,8 @@ help: the trait `TmaPolicy` is implemented for `AsyncCompletion` note: required by a bound in `TwoCtaRing::::producer_expect_tx` --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { -... | pub unsafe fn producer_expect_tx(self, ticket: u32, bytes: u32) + | ------------------ required by a bound in this associated function | where | Full: TmaPolicy, | ^^^^^^^^^ required by this bound in `TwoCtaRing::::producer_expect_tx` @@ -57,9 +55,8 @@ help: the trait `TmaPolicy` is implemented for `AsyncCompletion` note: required by a bound in `TwoCtaRing::::producer_expect_tx` --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { -... | pub unsafe fn producer_expect_tx(self, ticket: u32, bytes: u32) + | ------------------ required by a bound in this associated function | where | Full: TmaPolicy, | ^^^^^^^^^ required by this bound in `TwoCtaRing::::producer_expect_tx` @@ -87,9 +84,8 @@ help: the following other types implement trait `MmaPolicy` note: required by a bound in `TwoCtaRing::::mma_full` --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { -... | pub unsafe fn mma_full(self, ticket: u32) + | -------- required by a bound in this associated function | where | Full: MmaPolicy, | ^^^^^^^^^ required by this bound in `TwoCtaRing::::mma_full` @@ -117,9 +113,8 @@ help: the following other types implement trait `MmaPolicy` note: required by a bound in `TwoCtaRing::::mma_empty_cursor` --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { -... | pub unsafe fn mma_empty_cursor(self, cursor: RingCursor) + | ---------------- required by a bound in this associated function | where | Empty: MmaPolicy, | ^^^^^^^^^ required by this bound in `TwoCtaRing::::mma_empty_cursor` @@ -141,9 +136,8 @@ help: the following other types implement trait `SignalPolicy` note: required by a bound in `TwoCtaRing::::thread_full` --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { -... | pub unsafe fn thread_full(self, ticket: u32) + | ----------- required by a bound in this associated function | where | Full: SignalPolicy, | ^^^^^^^^^^^^ required by this bound in `TwoCtaRing::::thread_full` diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_wrong_signal.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_wrong_signal.stderr index f90258a26f..52fec6c5be 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_wrong_signal.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_wrong_signal.stderr @@ -12,9 +12,8 @@ help: the trait `TmaPolicy` is implemented for `AsyncCompletion` note: required by a bound in `TwoCtaRing::::producer_expect_tx` --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { -... | pub unsafe fn producer_expect_tx(self, ticket: u32, bytes: u32) + | ------------------ required by a bound in this associated function | where | Full: TmaPolicy, | ^^^^^^^^^ required by this bound in `TwoCtaRing::::producer_expect_tx` @@ -42,9 +41,8 @@ help: the following other types implement trait `MmaPolicy` note: required by a bound in `TwoCtaRing::::mma_full` --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { -... | pub unsafe fn mma_full(self, ticket: u32) + | -------- required by a bound in this associated function | where | Full: MmaPolicy, | ^^^^^^^^^ required by this bound in `TwoCtaRing::::mma_full` @@ -66,9 +64,8 @@ help: the following other types implement trait `SignalPolicy` note: required by a bound in `TwoCtaRing::::thread_empty` --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { -... | pub unsafe fn thread_empty(self, ticket: u32) + | ------------ required by a bound in this associated function | where | Empty: SignalPolicy, | ^^^^^^^^^^^^ required by this bound in `TwoCtaRing::::thread_empty` @@ -90,9 +87,8 @@ help: the following other types implement trait `SignalPolicy` note: required by a bound in `TwoCtaRing::::thread_full` --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { -... | pub unsafe fn thread_full(self, ticket: u32) + | ----------- required by a bound in this associated function | where | Full: SignalPolicy, | ^^^^^^^^^^^^ required by this bound in `TwoCtaRing::::thread_full` @@ -120,9 +116,8 @@ help: the following other types implement trait `MmaPolicy` note: required by a bound in `TwoCtaRing::::mma_empty` --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { -... | pub unsafe fn mma_empty(self, ticket: u32) + | --------- required by a bound in this associated function | where | Empty: MmaPolicy, | ^^^^^^^^^ required by this bound in `TwoCtaRing::::mma_empty` diff --git a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_zero_stages.stderr b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_zero_stages.stderr index 13c7ffaa97..7d26893413 100644 --- a/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_zero_stages.stderr +++ b/cuda-oxide/cute-rs-preview/cute-rs/tests/compile_fail/two_cta_ring_zero_stages.stderr @@ -1,8 +1,6 @@ error[E0080]: evaluation panicked: assertion failed: STAGES > 0 --> src/attention/sm100/pipeline.rs | - | impl TwoCtaRing { - | /// Reject zero-stage rings before modulo or storage-size calculations. | const VALID: () = assert!(STAGES > 0); | ^^^^^^^^^^^^^^^^^^^ evaluation of `cute_rs::attention::sm100::pipeline::TwoCtaRing::<0, cute_rs::attention::sm100::pipeline::AsyncCompletion>::VALID` failed here diff --git a/cuda-oxide/intrinsics/upstream.lock b/cuda-oxide/intrinsics/upstream.lock index e13c5668f9..36d21387d4 100644 --- a/cuda-oxide/intrinsics/upstream.lock +++ b/cuda-oxide/intrinsics/upstream.lock @@ -7,7 +7,7 @@ provenance = "Public rust-lang/llvm-project source revision used by Rust commit public_output_allowed = true [rust_toolchain] -commit_hash = "6f5da406b56c2db6607202b63daee73aa8c61297" +commit_hash = "58cf8f9c35b355199727627d07ae34a4e2846d67" [llvm_tblgen] name = "llvm-tblgen" diff --git a/cuda-oxide/rust-toolchain.toml b/cuda-oxide/rust-toolchain.toml index 1f9d49c31d..c827fac49c 100644 --- a/cuda-oxide/rust-toolchain.toml +++ b/cuda-oxide/rust-toolchain.toml @@ -1,5 +1,5 @@ [toolchain] -channel = "nightly-2026-10-11" +channel = "1.99.1" components = [ "rust-src", "rustc-dev", diff --git a/cuda-oxide/scripts/check-dependency-licenses.sh b/cuda-oxide/scripts/check-dependency-licenses.sh index 3074df0a02..778581f439 100755 --- a/cuda-oxide/scripts/check-dependency-licenses.sh +++ b/cuda-oxide/scripts/check-dependency-licenses.sh @@ -25,6 +25,7 @@ set -euo pipefail # dictionary collation, reject the pair rustc-hash/rustc_apfloat with # "input is not in sorted order", and abort the run via set -e. export LC_ALL=C +export RUSTC_BOOTSTRAP=1 cd "$(dirname "$0")/.." GIT_ROOT="$(git rev-parse --show-toplevel)" diff --git a/cuda-oxide/scripts/check-example-license-policy.sh b/cuda-oxide/scripts/check-example-license-policy.sh index 53adab9e55..29859438ce 100644 --- a/cuda-oxide/scripts/check-example-license-policy.sh +++ b/cuda-oxide/scripts/check-example-license-policy.sh @@ -33,6 +33,7 @@ set -euo pipefail export LC_ALL=C +export RUSTC_BOOTSTRAP=1 cd "$(dirname "$0")/.." diff --git a/cuda-oxide/scripts/check-primitives-api-compile.py b/cuda-oxide/scripts/check-primitives-api-compile.py index 2b0beb6ac1..bee4ac54d8 100644 --- a/cuda-oxide/scripts/check-primitives-api-compile.py +++ b/cuda-oxide/scripts/check-primitives-api-compile.py @@ -345,7 +345,7 @@ def main(): run_dir = Path(tempfile.mkdtemp(prefix="run-", dir=evidence)) env = os.environ.copy() channel = tomllib.loads((ROOT / "rust-toolchain.toml").read_text())["toolchain"]["channel"] - env.update(RUSTUP_TOOLCHAIN=channel, CARGO_ENCODED_RUSTFLAGS="", RUSTFLAGS="") + env.update(RUSTUP_TOOLCHAIN=channel, RUSTC_BOOTSTRAP="1", CARGO_ENCODED_RUSTFLAGS="", RUSTFLAGS="") sysroot = subprocess.check_output(["rustc", "--print", "sysroot"], cwd=ROOT, env=env, text=True).strip() rustc_version = subprocess.check_output(["rustc", "-vV"], cwd=ROOT, env=env, text=True) host = re.search(r"^host: (.+)$", rustc_version, re.M).group(1) diff --git a/cuda-oxide/scripts/check-toolchain-parity.sh b/cuda-oxide/scripts/check-toolchain-parity.sh index bc31be6bf5..099c7df9f2 100755 --- a/cuda-oxide/scripts/check-toolchain-parity.sh +++ b/cuda-oxide/scripts/check-toolchain-parity.sh @@ -236,31 +236,28 @@ if not quoted: "guard) or the fence style changed (fix the pattern)" ) -# Every dated nightly reference in tracked markdown. The book's install -# pages and the READMEs spell the pin inside commands a reader runs outside -# a checkout (`rustup toolchain install nightly-...`, -# `cargo +nightly-... install`), where no rust-toolchain.toml can correct a -# stale date, plus prose naming the pin. The token itself is the check -- -# any dated nightly a tracked page mentions must be the pinned one -- so the -# guard survives pages being reworded, moved, or added. Component lists on -# `rustup component add` lines stay uncovered, as above; the channel date on -# those same lines is checked like any other. -DATED = re.compile(r"nightly-\d{4}-\d{2}-\d{2}") -if not DATED.fullmatch(root_channel): +# Every explicit toolchain reference in tracked markdown (excluding historical +# intrinsic capture evidence under `cuda-oxide/intrinsics/`). The book's +# install pages and the READMEs spell the pin inside commands a reader runs +# outside a checkout (`rustup toolchain install `, `cargo + install`, +# `rustup component add ... --toolchain `), where no rust-toolchain.toml +# can correct a stale version, plus prose naming the pin. +EXACT_PIN = re.compile(r"(?:nightly-\d{4}-\d{2}-\d{2}|\d+\.\d+\.\d+)") +if not EXACT_PIN.fullmatch(root_channel): sys.exit( f"parse self-test failed: {root_path} channel {root_channel!r} is " - "not a dated nightly; this guard assumes an exact pin" + "not an exact release or dated nightly pin" ) markdown = sorted( path for path in subprocess.run( - ["git", "-C", git_root, "ls-files", "-z", "--", "*.md"], + ["git", "-C", git_root, "ls-files", "-z", "--", "cuda-oxide/*.md", "cuda-oxide/**/*.md"], stdout=subprocess.PIPE, check=True, ) .stdout.decode() .split("\0") - if path + if path and not path.startswith("cuda-oxide/intrinsics/") ) if len(markdown) < 20: sys.exit( @@ -268,20 +265,33 @@ if len(markdown) < 20: "markdown files" ) -dated = 0 +DATED_NIGHTLY = re.compile(r"nightly-\d{4}-\d{2}-\d{2}") +COMMAND_PIN = re.compile( + r"(?:rustup\s+toolchain\s+install\s+|" + r"--toolchain\s+|" + r"cargo\s+\+)([A-Za-z0-9._-]+)" +) +pinned_refs = 0 for path in markdown: for number, line in enumerate(read(os.path.join(git_root, path)).splitlines(), start=1): - for token in DATED.findall(line): - dated += 1 + for token in DATED_NIGHTLY.findall(line): + pinned_refs += 1 + if token != root_channel: + failures.append( + f"{path}:{number} spells {token!r}, " + f"{root_path} pins {root_channel!r}" + ) + for token in COMMAND_PIN.findall(line): + pinned_refs += 1 if token != root_channel: failures.append( f"{path}:{number} spells {token!r}, " f"{root_path} pins {root_channel!r}" ) -if not dated: +if not pinned_refs: sys.exit( - "parse self-test failed: no dated nightly reference found in any " + "parse self-test failed: no toolchain command reference found in any " "tracked markdown; either the docs stopped spelling the pin (delete " "this block of the guard) or the pattern is stale (fix it)" ) @@ -295,7 +305,7 @@ if failures: f"Every copy must state the same channel and components as {root_path}. " "The `cargo oxide new` scaffold must match exactly " "(rustc_private dylibs);\nthe book's blocks are presented to readers as " - "the repo's real file, and the\ndated commands run where no " + "the repo's real file, and the\nexplicit toolchain commands run where no " "rust-toolchain.toml can correct them.", file=sys.stderr, ) @@ -304,7 +314,7 @@ if failures: print( f"OK: {root_path} pins {root_channel} with {len(root_components)} components; " "the `cargo oxide new` scaffold, the devcontainer, all " - f"{quoted} block(s) quoted in the book, and all {dated} dated reference(s) " + f"{quoted} block(s) quoted in the book, and all {pinned_refs} toolchain reference(s) " "in tracked markdown agree." ) PY