From e9cd672f05e65729af3f37040c312265a1ea9d26 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Fri, 12 Jun 2026 15:37:01 +0530 Subject: [PATCH 1/6] refactor(cuda-host): move libdevice discovery into libnvvm-sys libdevice ships in the toolkit's nvvm/ component next to libnvvm.so and is consumed together with libNVVM in the LTOIR pipeline, so the probe belongs in libnvvm-sys. cuda_host::ltoir::find_libdevice becomes a thin delegating wrapper with unchanged error type and message text, and ltoir.rs drops its now-unused cuda_roots/cuda_roots_from_env copy (the same helper already lives in libnvvm-sys). This also lets cargo-oxide's doctor probe libdevice without pulling the cuda-host -> cuda-core -> cuda-bindings chain. --- crates/cuda-host/src/ltoir.rs | 64 ++------------------ crates/libnvvm-sys/src/lib.rs | 111 +++++++++++++++++++++++++++++++++- 2 files changed, 115 insertions(+), 60 deletions(-) diff --git a/crates/cuda-host/src/ltoir.rs b/crates/cuda-host/src/ltoir.rs index 0518a461a8..a005ebefb8 100644 --- a/crates/cuda-host/src/ltoir.rs +++ b/crates/cuda-host/src/ltoir.rs @@ -296,24 +296,12 @@ pub fn load_kernel_module( /// /// Returns [`LtoirError::LibdeviceNotFound`] with the full list of probed /// paths if nothing matches. +/// +/// Thin wrapper over [`libnvvm_sys::find_libdevice`], which owns the probe +/// (libdevice ships in the toolkit's `nvvm/` component next to `libnvvm.so`). pub fn find_libdevice() -> Result { - if let Ok(p) = std::env::var("CUDA_OXIDE_LIBDEVICE") { - let path = PathBuf::from(p); - if path.exists() { - return Ok(path); - } - } - let mut tried = Vec::new(); - for root in cuda_roots() { - let candidate = root.join("nvvm/libdevice/libdevice.10.bc"); - tried.push(candidate.display().to_string()); - if candidate.exists() { - return Ok(candidate); - } - } - Err(LtoirError::LibdeviceNotFound { - tried: tried.join("\n "), - }) + libnvvm_sys::find_libdevice() + .map_err(|libnvvm_sys::LibdeviceNotFound { tried }| LtoirError::LibdeviceNotFound { tried }) } /// Read the GPU arch (`sm_XX`) for the cubin build, defaulting to `sm_120` @@ -352,22 +340,6 @@ fn manifest_dir() -> PathBuf { // Internal utilities // ============================================================================ -fn cuda_roots() -> Vec { - cuda_roots_from_env(|var| std::env::var(var).ok()) -} - -fn cuda_roots_from_env(mut get_env: impl FnMut(&str) -> Option) -> Vec { - let mut roots = Vec::new(); - for var in ["CUDA_TOOLKIT_PATH", "CUDA_HOME", "CUDA_PATH"] { - if let Some(r) = get_env(var) { - roots.push(PathBuf::from(r)); - } - } - roots.push(PathBuf::from("/usr/local/cuda")); - roots.push(PathBuf::from("/opt/cuda")); - roots -} - /// Convert `sm_120` to `compute_120`. Returns the input unchanged if it /// doesn't start with `sm_`. fn sm_to_compute(arch: &str) -> String { @@ -396,29 +368,3 @@ fn needs_rebuild(target: &Path, sources: &[&Path]) -> bool { } false } - -#[cfg(test)] -mod tests { - use super::*; - - #[test] - fn cuda_roots_prefers_project_toolkit_env_var() { - let roots = cuda_roots_from_env(|var| match var { - "CUDA_TOOLKIT_PATH" => Some("/cuda/toolkit".to_string()), - "CUDA_HOME" => Some("/cuda/home".to_string()), - "CUDA_PATH" => Some("/cuda/path".to_string()), - _ => None, - }); - - assert_eq!( - roots, - vec![ - PathBuf::from("/cuda/toolkit"), - PathBuf::from("/cuda/home"), - PathBuf::from("/cuda/path"), - PathBuf::from("/usr/local/cuda"), - PathBuf::from("/opt/cuda"), - ] - ); - } -} diff --git a/crates/libnvvm-sys/src/lib.rs b/crates/libnvvm-sys/src/lib.rs index 888f342719..691fedc977 100644 --- a/crates/libnvvm-sys/src/lib.rs +++ b/crates/libnvvm-sys/src/lib.rs @@ -41,7 +41,7 @@ use libloading::{Library, Symbol}; use std::ffi::{CString, c_char, c_int, c_void}; -use std::path::PathBuf; +use std::path::{Path, PathBuf}; use std::ptr; use thiserror::Error; @@ -115,6 +115,17 @@ pub enum NvvmError { }, } +/// `libdevice.10.bc` could not be located on this system. `tried` lists +/// every path that was probed, in order, joined by newlines. +#[derive(Debug, Error)] +#[error( + "Could not locate libdevice.10.bc. Set CUDA_OXIDE_LIBDEVICE, CUDA_TOOLKIT_PATH, or CUDA_HOME, or install the CUDA Toolkit. Tried:\n {tried}" +)] +pub struct LibdeviceNotFound { + /// Newline-joined list of paths that were probed. + pub tried: String, +} + // ============================================================================ // Library handle // ============================================================================ @@ -416,6 +427,52 @@ fn cuda_roots_from_env(mut get_env: impl FnMut(&str) -> Option) -> Vec

/nvvm/libdevice/libdevice.10.bc` for `` in +/// `CUDA_TOOLKIT_PATH`, `CUDA_HOME`, `CUDA_PATH`, `/usr/local/cuda`, +/// `/opt/cuda`. +/// +/// Returns [`LibdeviceNotFound`] with the full list of probed paths if +/// nothing matches. +pub fn find_libdevice() -> Result { + find_libdevice_with(|var| std::env::var(var).ok(), |path| path.exists()) +} + +fn find_libdevice_with( + mut get_env: impl FnMut(&str) -> Option, + mut exists: impl FnMut(&Path) -> bool, +) -> Result { + if let Some(p) = get_env("CUDA_OXIDE_LIBDEVICE") { + let path = PathBuf::from(p); + if exists(&path) { + return Ok(path); + } + } + let mut tried = Vec::new(); + for root in cuda_roots_from_env(&mut get_env) { + let candidate = root.join("nvvm/libdevice/libdevice.10.bc"); + tried.push(candidate.display().to_string()); + if exists(&candidate) { + return Ok(candidate); + } + } + Err(LibdeviceNotFound { + tried: tried.join("\n "), + }) +} + #[cfg(test)] mod tests { use super::*; @@ -440,4 +497,56 @@ mod tests { ] ); } + + #[test] + fn find_libdevice_honors_explicit_override_file() { + let found = find_libdevice_with( + |var| (var == "CUDA_OXIDE_LIBDEVICE").then(|| "/elsewhere/libdevice.10.bc".to_string()), + |path| path == Path::new("/elsewhere/libdevice.10.bc"), + ); + + assert_eq!(found.unwrap(), PathBuf::from("/elsewhere/libdevice.10.bc")); + } + + #[test] + fn find_libdevice_probes_roots_in_order() { + // CUDA_HOME has the file, but CUDA_TOOLKIT_PATH is probed first and + // also has it; the first match must win. + let found = find_libdevice_with( + |var| match var { + "CUDA_TOOLKIT_PATH" => Some("/cuda/toolkit".to_string()), + "CUDA_HOME" => Some("/cuda/home".to_string()), + _ => None, + }, + |path| { + path == Path::new("/cuda/toolkit/nvvm/libdevice/libdevice.10.bc") + || path == Path::new("/cuda/home/nvvm/libdevice/libdevice.10.bc") + }, + ); + + assert_eq!( + found.unwrap(), + PathBuf::from("/cuda/toolkit/nvvm/libdevice/libdevice.10.bc") + ); + } + + #[test] + fn find_libdevice_failure_lists_every_probed_path() { + let err = find_libdevice_with( + |var| (var == "CUDA_HOME").then(|| "/cuda/home".to_string()), + |_| false, + ) + .unwrap_err(); + + assert_eq!( + err.tried, + "/cuda/home/nvvm/libdevice/libdevice.10.bc\n \ + /usr/local/cuda/nvvm/libdevice/libdevice.10.bc\n \ + /opt/cuda/nvvm/libdevice/libdevice.10.bc" + ); + let message = err.to_string(); + assert!(message.contains("CUDA_OXIDE_LIBDEVICE")); + assert!(message.contains("CUDA_TOOLKIT_PATH")); + assert!(message.contains("CUDA_HOME")); + } } From c1d28d795518752d0922f7708458b42cbb2fb651 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Fri, 12 Jun 2026 15:39:44 +0530 Subject: [PATCH 2/6] feat(cargo-oxide): drop cuda-core/cuda-host so the subcommand builds without the CUDA toolkit cargo-oxide had exactly two real CUDA-coupled call sites: the doctor libdevice probe (cuda_host::ltoir::find_libdevice) and the run-target auto-detect (cuda_core::CudaContext compute_capability). Both pulled the cuda-host -> cuda-core -> cuda-bindings chain, whose build script hard-fails without cuda.h and links libcuda, so on a bare machine 'cargo oxide doctor' died in a build script before main() (issue #87) and a prebuilt binary could not even load without a libcuda.so.1 stub. - doctor's libdevice check now calls libnvvm_sys::find_libdevice() (same probe, no cuda-bindings in the graph) - detect_run_target_arch now derives the compute capability from 'nvidia-smi --query-gpu=compute_cap --format=csv,noheader', the same derivation scripts/smoketest.sh already uses; output is trusted only on exit 0 AND a . shaped first line, since nvidia-smi prints failure banners to stdout - precedence is unchanged: --arch and CUDA_OXIDE_TARGET still override, the detected arch stays an advisory CUDA_OXIDE_DEVICE_ARCH hint, and format_sm_arch still emits sm_XYa for cc >= 9.0 Known tradeoff, documented on the helper: nvidia-smi lists GPUs in PCI order while CUDA defaults to fastest-first, so on heterogeneous multi-GPU boxes the hint may describe a different GPU than CUDA device 0. Safe because the hint is advisory and hard overrides remain. cargo tree -p cargo-oxide -e normal,build now contains no cuda-bindings, cuda-core, cuda-host, bindgen, or clang-sys, and the binary's DT_NEEDED no longer carries libcuda.so.1. --- Cargo.lock | 2 - crates/cargo-oxide/Cargo.toml | 7 +- crates/cargo-oxide/src/commands.rs | 126 +++++++++++++++++++++++------ 3 files changed, 108 insertions(+), 27 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index 4e3e909c3d..44f226dae9 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -191,8 +191,6 @@ name = "cargo-oxide" version = "0.2.1" dependencies = [ "clap", - "cuda-core", - "cuda-host", "libnvvm-sys", "nvjitlink-sys", "toml", diff --git a/crates/cargo-oxide/Cargo.toml b/crates/cargo-oxide/Cargo.toml index c3afd59b5e..8f3c88ac4a 100644 --- a/crates/cargo-oxide/Cargo.toml +++ b/crates/cargo-oxide/Cargo.toml @@ -11,10 +11,13 @@ description = "Cargo subcommand for cuda-oxide: build and run Rust GPU programs" name = "cargo-oxide" path = "src/main.rs" +# Deliberately CUDA-free: no dependency here may pull `cuda-bindings` +# (or any other build script that needs the CUDA toolkit), so that +# `cargo oxide doctor` builds and runs on a machine with no toolkit and +# no driver. libnvvm-sys / nvjitlink-sys are pure-Rust runtime-dlopen +# crates and are fine. [dependencies] clap = { version = "4", features = ["derive"] } libnvvm-sys = { workspace = true } nvjitlink-sys = { workspace = true } -cuda-host = { workspace = true } -cuda-core = { workspace = true } toml = "1" diff --git a/crates/cargo-oxide/src/commands.rs b/crates/cargo-oxide/src/commands.rs index 3a94cab7c4..ce98b19b33 100644 --- a/crates/cargo-oxide/src/commands.rs +++ b/crates/cargo-oxide/src/commands.rs @@ -109,7 +109,7 @@ pub fn codegen_run( // Target precedence for `cargo oxide run` (highest first): // 1. --arch explicit user override -> CUDA_OXIDE_TARGET // 2. CUDA_OXIDE_TARGET= explicit env override (from the parent) - // 3. detected GPU arch of CUDA device 0 -> CUDA_OXIDE_DEVICE_ARCH (a hint) + // 3. detected GPU arch (via nvidia-smi) -> CUDA_OXIDE_DEVICE_ARCH (a hint) // 4. backend feature-based default (`select_target` in mir-importer) // // Slot 3 is a HINT, not an override: the backend builds for the detected @@ -117,8 +117,8 @@ pub fn codegen_run( // arch (tcgen05 needs sm_100a even on a consumer sm_120 GPU), the backend // builds for the required arch and the module simply skips at load time. // We only detect for `run`, not `build`/`pipeline`: `run` loads the cubin - // on device 0, whereas those may legitimately cross-compile for another - // machine. + // on the local GPU, whereas those may legitimately cross-compile for + // another machine. let detected_device_arch = detect_run_target_arch(arch, emit_nvvm_ir); if let Some(interop) = interop.filter(|config| !config.device_crates.is_empty()) { @@ -155,7 +155,7 @@ pub fn codegen_run( // a hard target: the backend builds for it unless a kernel needs a // newer arch (e.g. tcgen05 forces sm_100a even on a consumer sm_120 // GPU), so the final PTX target may differ. - println!("Detected GPU arch: {dev} (CUDA device 0)"); + println!("Detected GPU arch: {dev} (via nvidia-smi)"); println!(); } println!("This is the proper cargo workflow:"); @@ -244,7 +244,7 @@ fn codegen_run_interop( println!("Interop kind: {}", kind); } if let Some(dev) = detected_device_arch { - println!("Detected GPU arch: {dev} (CUDA device 0)"); + println!("Detected GPU arch: {dev} (via nvidia-smi)"); } println!(); @@ -912,7 +912,7 @@ pub fn doctor(ctx: &Context) { } print!("libdevice (libdevice.10.bc)... "); - match cuda_host::ltoir::find_libdevice() { + match libnvvm_sys::find_libdevice() { Ok(path) => println!("✓ {}", path.display()), Err(e) => { println!("✗ {}", e); @@ -1344,11 +1344,11 @@ fn apply_device_arch_hint( /// 1. `--arch ` (explicit user override) /// 2. `CUDA_OXIDE_TARGET=` (explicit env override, set in the parent /// process before invoking `cargo oxide run`) -/// 3. **This function**: the compute capability of the GPU in CUDA device 0, -/// forwarded as the `CUDA_OXIDE_DEVICE_ARCH` *hint*. Emits the arch-specific -/// `sm_XYa` form for cc >= 9.0 (so the backend can lower WGMMA / tcgen05 / -/// TMA-multicast when the GPU supports them) and the plain `sm_XY` form for -/// cc < 9.0. +/// 3. **This function**: the compute capability of the first GPU reported by +/// `nvidia-smi`, forwarded as the `CUDA_OXIDE_DEVICE_ARCH` *hint*. Emits +/// the arch-specific `sm_XYa` form for cc >= 9.0 (so the backend can lower +/// WGMMA / tcgen05 / TMA-multicast when the GPU supports them) and the +/// plain `sm_XY` form for cc < 9.0. /// 4. Backend feature-based default (`select_target` in /// `mir-importer::pipeline`), which picks the minimum `sm_XX` required by /// the IR shape (e.g. `Basic -> sm_80`, `Cluster -> sm_90`, `Tma -> sm_100`). @@ -1360,11 +1360,11 @@ fn apply_device_arch_hint( /// /// # Why only `run` /// -/// `run` immediately loads the generated module on device 0 and launches the -/// kernel, so a target older than the local GPU's compute capability is the -/// only safe default. `build` and `pipeline` may legitimately cross-compile -/// to a different machine, so they keep the backend's feature-based default -/// untouched. +/// `run` immediately loads the generated module on the local GPU and launches +/// the kernel, so a target older than the local GPU's compute capability is +/// the only safe default. `build` and `pipeline` may legitimately +/// cross-compile to a different machine, so they keep the backend's +/// feature-based default untouched. /// /// # Why this is needed even with the backend default /// @@ -1380,18 +1380,61 @@ fn apply_device_arch_hint( /// - `CUDA_OXIDE_TARGET` is set in the environment (slot 2 wins). /// - `--emit-nvvm-ir` is in effect (NVVM IR mode requires explicit `--arch`, /// enforced by the CLI parser). -/// - No CUDA driver / device 0 is available on the machine (CI runners without -/// GPUs, headless build boxes). The caller falls through to slot 4 and -/// the backend's feature-based default applies. +/// - No CUDA driver / GPU is available on the machine (CI runners without +/// GPUs, headless build boxes), or `nvidia-smi` is missing or broken. The +/// caller falls through to slot 4 and the backend's feature-based default +/// applies. fn detect_run_target_arch(arch: Option<&str>, emit_nvvm_ir: bool) -> Option { if arch.is_some() || emit_nvvm_ir || std::env::var_os("CUDA_OXIDE_TARGET").is_some() { return None; } - cuda_core::CudaContext::new(0) - .and_then(|ctx| ctx.compute_capability()) + query_device_compute_cap().map(format_sm_arch) +} + +/// Query the compute capability of the first GPU via `nvidia-smi`. +/// +/// Runs `nvidia-smi --query-gpu=compute_cap --format=csv,noheader` and parses +/// the first output line. A subprocess probe (rather than the CUDA driver +/// API) keeps cargo-oxide free of any link-time or dlopen dependency on +/// `libcuda`, so the subcommand builds and runs on machines with no CUDA +/// toolkit and no driver; `scripts/smoketest.sh` derives `sm_XX` from +/// `nvidia-smi` the same way. +/// +/// Caveat: `nvidia-smi` enumerates GPUs in PCI bus order, while CUDA's +/// default device order is fastest-first, so on heterogeneous multi-GPU +/// machines this may describe a different GPU than CUDA device 0. That is +/// safe because `CUDA_OXIDE_DEVICE_ARCH` is advisory (the backend only +/// honors a compatible hint) and `--arch` / `CUDA_OXIDE_TARGET` remain hard +/// overrides. +fn query_device_compute_cap() -> Option<(u32, u32)> { + let output = Command::new("nvidia-smi") + .args(["--query-gpu=compute_cap", "--format=csv,noheader"]) + .output() .ok() - .map(format_sm_arch) + .filter(|o| o.status.success())?; + parse_compute_cap(&String::from_utf8_lossy(&output.stdout)) +} + +/// Parse the first line of `nvidia-smi --query-gpu=compute_cap` output as a +/// `(major, minor)` compute-capability pair. Returns `None` for anything +/// that is not shaped `.`. +fn parse_compute_cap(stdout: &str) -> Option<(u32, u32)> { + parse_compute_cap_field(stdout.lines().next()?) +} + +/// Parse a single `compute_cap` CSV field (e.g. `"12.0"`). +/// +/// Only the `.` shape is accepted: `nvidia-smi` prints its +/// failure banners ("NVIDIA-SMI has failed ...") to *stdout*, sometimes with +/// exit status 0, so this shape check is the real gate, not the exit status. +fn parse_compute_cap_field(field: &str) -> Option<(u32, u32)> { + let (major, minor) = field.trim().split_once('.')?; + let all_digits = |s: &str| !s.is_empty() && s.bytes().all(|b| b.is_ascii_digit()); + if !all_digits(major) || !all_digits(minor) { + return None; + } + Some((major.parse().ok()?, minor.parse().ok()?)) } /// Format a `(major, minor)` compute-capability tuple as the `sm_XX` / @@ -1421,7 +1464,7 @@ fn detect_run_target_arch(arch: Option<&str>, emit_nvvm_ir: bool) -> Option String { +fn format_sm_arch((major, minor): (u32, u32)) -> String { if major >= 9 { format!("sm_{}{}a", major, minor) } else { @@ -1910,6 +1953,43 @@ mod tests { assert_eq!(format_sm_arch((12, 0)), "sm_120a"); // consumer Blackwell } + #[test] + fn parse_compute_cap_accepts_real_nvidia_smi_output() { + assert_eq!(parse_compute_cap("12.0\n"), Some((12, 0))); + assert_eq!(parse_compute_cap("7.5\n"), Some((7, 5))); + assert_eq!(parse_compute_cap("10.3"), Some((10, 3))); + // End-to-end with format_sm_arch: the values the backend sees. + assert_eq!( + format_sm_arch(parse_compute_cap("12.0\n").unwrap()), + "sm_120a" + ); + assert_eq!(format_sm_arch(parse_compute_cap("7.5\n").unwrap()), "sm_75"); + } + + #[test] + fn parse_compute_cap_takes_first_gpu_on_multi_gpu_machines() { + assert_eq!(parse_compute_cap("9.0\n12.0\n"), Some((9, 0))); + } + + #[test] + fn parse_compute_cap_rejects_failure_banners_and_garbage() { + // nvidia-smi prints failure text to STDOUT, not stderr. + assert_eq!( + parse_compute_cap( + "NVIDIA-SMI has failed because it couldn't communicate \ + with the NVIDIA driver.\n" + ), + None + ); + assert_eq!(parse_compute_cap(""), None); + assert_eq!(parse_compute_cap("\n"), None); + assert_eq!(parse_compute_cap("N/A\n"), None); + assert_eq!(parse_compute_cap("12\n"), None); + assert_eq!(parse_compute_cap("12.\n"), None); + assert_eq!(parse_compute_cap(".5\n"), None); + assert_eq!(parse_compute_cap("12.0.1\n"), None); + } + #[test] fn detect_run_target_arch_skips_when_arch_explicit() { // --arch wins; never query the GPU. From 825ab20f3ca3010030727bc77defb084a33cf109 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Fri, 12 Jun 2026 15:45:27 +0530 Subject: [PATCH 3/6] feat(cargo-oxide): teach doctor to diagnose missing cuda.h and absent driver Fixes #87. With the cuda-bindings edge gone, doctor finally runs on a bare machine, so make it actually diagnose one: - new fatal 'CUDA headers (cuda.h)' check that mirrors the discovery in cuda-bindings/build.rs (CUDA_TOOLKIT_PATH > CUDA_HOME > /usr/local/cuda; include/ then targets// include per host arch). On failure it lists every probed path and the env vars to set. The probe is kept in lockstep by hand (build.rs logic is not a library and cuda-bindings must stay out of cargo-oxide's graph); both helpers are pure and unit-tested - new informational 'NVIDIA driver / GPU' check via nvidia-smi (reports name + compute capability), with a /proc/driver/nvidia/ version fallback that distinguishes 'driver loaded but nvidia-smi broken' from 'no driver at all'. Non-fatal: only 'cargo oxide run' needs a driver, and GPU-less CI is a supported workflow - doctor is now build-free: the Doctor arm resolves a side-effect-free context (new resolve_doctor_context + backend::backend_so_candidate) instead of resolve_context, which unconditionally built the backend (and could git-clone in standalone mode) before printing anything - the backend-.so check downgrades to informational '- not built yet (run `cargo oxide setup`)': run/build/pipeline build it on demand, so a fresh healthy clone stays green README and book updated to describe the new checks and that cargo-oxide plus doctor work without the CUDA toolkit. --- crates/cargo-oxide/README.md | 12 +- crates/cargo-oxide/src/backend.rs | 28 +++ crates/cargo-oxide/src/commands.rs | 235 +++++++++++++++++- crates/cargo-oxide/src/main.rs | 4 +- .../getting-started/installation.md | 11 +- 5 files changed, 277 insertions(+), 13 deletions(-) diff --git a/crates/cargo-oxide/README.md b/crates/cargo-oxide/README.md index f5d14bcfb5..e1c58d4162 100644 --- a/crates/cargo-oxide/README.md +++ b/crates/cargo-oxide/README.md @@ -113,7 +113,17 @@ Formats all crates in the workspace: root workspace, `rustc-codegen-cuda`, and a ### `cargo oxide doctor` -Validates that your environment is correctly set up: Rust nightly toolchain, CUDA toolkit (`nvcc`), LLVM (`llc`), and the codegen backend `.so`. +Validates that your environment is correctly set up: Rust nightly toolchain, +CUDA headers (`cuda.h`), CUDA toolkit (`nvcc`, libNVVM, nvJitLink, +libdevice), LLVM (`llc`), clang/libclang, the NVIDIA driver / GPU, and the +codegen backend `.so`. Every check reports what was found or how to fix it. + +`cargo-oxide` itself builds and runs without the CUDA toolkit and without an +NVIDIA driver, and `doctor` never builds anything first, so it works on a +bare machine and tells you exactly what is missing. The driver / GPU check is +informational (only `cargo oxide run` needs a GPU), and a missing backend +`.so` just points at `cargo oxide setup` (`run`/`build` build it on demand +anyway). ### `cargo oxide setup` diff --git a/crates/cargo-oxide/src/backend.rs b/crates/cargo-oxide/src/backend.rs index 25a29a8522..a67b8b22a9 100644 --- a/crates/cargo-oxide/src/backend.rs +++ b/crates/cargo-oxide/src/backend.rs @@ -85,6 +85,34 @@ pub fn find_or_build_backend(workspace_root: &Path) -> PathBuf { auto_fetch_and_build() } +/// Returns where the backend `.so` lives (or would live), with NO side +/// effects: never builds, never clones, never touches the network. +/// +/// Mirrors the discovery order of [`find_or_build_backend`] minus its +/// build/clone steps: +/// +/// 1. `CUDA_OXIDE_BACKEND` env var, returned even when the file is missing +/// so the caller can report the configured-but-absent path. +/// 2. Local repo build path (`crates/rustc-codegen-cuda/target/debug/...`). +/// 3. Cache path at `~/.cargo/cuda-oxide/librustc_codegen_cuda.so`. +/// +/// `cargo oxide doctor` uses this so that a diagnostic run never triggers a +/// multi-minute backend build or a git clone before it can print anything. +pub fn backend_so_candidate(workspace_root: &Path) -> PathBuf { + if let Ok(path) = std::env::var("CUDA_OXIDE_BACKEND") { + return PathBuf::from(path); + } + + let codegen_crate = workspace_root.join("crates/rustc-codegen-cuda"); + if codegen_crate.is_dir() { + return codegen_crate.join("target/debug/librustc_codegen_cuda.so"); + } + + cache_directory() + .map(|dir| dir.join("librustc_codegen_cuda.so")) + .unwrap_or_else(|| PathBuf::from("librustc_codegen_cuda.so")) +} + /// Returns true when the cached backend `.so` is older than the running /// `cargo-oxide` binary, which means the user has upgraded the binary /// since the cache was last built. diff --git a/crates/cargo-oxide/src/commands.rs b/crates/cargo-oxide/src/commands.rs index ce98b19b33..e5b5d26339 100644 --- a/crates/cargo-oxide/src/commands.rs +++ b/crates/cargo-oxide/src/commands.rs @@ -78,6 +78,51 @@ pub fn resolve_context() -> Context { std::process::exit(1); } +/// Resolve a context for `cargo oxide doctor` with NO side effects. +/// +/// Identical discovery to [`resolve_context`], except the backend `.so` is +/// only *located* (via [`backend::backend_so_candidate`]), never built and +/// never cloned. A diagnostic command must be runnable on a machine where +/// nothing is set up yet; gating it behind a multi-minute backend build (or +/// a network clone) would hide the very problems it exists to report. +/// `run`/`build`/`pipeline`/`setup` still build the backend on demand. +pub fn resolve_doctor_context() -> Context { + if let Some(workspace_root) = backend::find_workspace_root() { + let codegen_crate = workspace_root.join("crates/rustc-codegen-cuda"); + let examples_dir = codegen_crate.join("examples"); + let backend_so = backend::backend_so_candidate(&workspace_root); + return Context { + workspace_root, + codegen_crate, + examples_dir, + backend_so, + is_workspace: true, + }; + } + + let cwd = std::env::current_dir().unwrap_or_else(|e| { + eprintln!("Error: cannot determine current directory: {}", e); + std::process::exit(1); + }); + + if cwd.join("Cargo.toml").is_file() { + let backend_so = backend::backend_so_candidate(&cwd); + return Context { + workspace_root: cwd.clone(), + codegen_crate: cwd.clone(), + examples_dir: cwd.clone(), + backend_so, + is_workspace: false, + }; + } + + eprintln!("Error: Could not find cuda-oxide workspace or a standalone Cargo.toml."); + eprintln!(); + eprintln!("Run from inside the cuda-oxide repository, or from a project created"); + eprintln!("with `cargo oxide new `."); + std::process::exit(1); +} + // ============================================================================= // Run command // ============================================================================= @@ -814,8 +859,15 @@ fn run_cargo_fmt(dir: &Path, check: bool) -> bool { /// Validate the development environment. /// /// Checks for: Rust nightly toolchain, `rust-toolchain.toml`, the codegen -/// backend `.so`, CUDA toolkit (`nvcc`), LLVM (`llc`), and optionally -/// `cuda-gdb`. Exits non-zero if any required check fails. +/// backend `.so` (informational), CUDA headers (`cuda.h`), CUDA toolkit +/// (`nvcc`, libNVVM, nvJitLink, libdevice), LLVM (`llc`), clang/libclang, +/// the NVIDIA driver / GPU (informational), and optionally `cuda-gdb`. +/// Exits non-zero if any required check fails. +/// +/// Doctor itself needs neither the CUDA toolkit nor a driver: every check +/// is a subprocess, a filesystem probe, or a runtime `dlopen`, and the +/// caller resolves the context via [`resolve_doctor_context`] so nothing is +/// built first. This is what lets it diagnose a bare machine (issue #87). pub fn doctor(ctx: &Context) { println!("cargo-oxide environment check"); println!("=============================="); @@ -852,16 +904,39 @@ pub fn doctor(ctx: &Context) { ok = false; } - // 3. Backend .so + // 3. Backend .so. Informational, not fatal: `run`/`build`/`pipeline` + // build the backend on demand, so "not built yet" is a healthy state + // for a fresh clone. print!("Codegen backend... "); if ctx.backend_so.exists() { println!("✓ {}", ctx.backend_so.display()); } else { - println!("✗ not found (run `cargo oxide setup`)"); - ok = false; + println!("- not built yet (run `cargo oxide setup`)"); + } + + // 4. CUDA headers (cuda.h). The host `cuda-bindings` crate cannot build + // without them; cargo-oxide itself deliberately can, which is what makes + // this check reachable on a toolkit-less machine instead of dying inside + // cuda-bindings' build script (issue #87). + print!("CUDA headers (cuda.h)... "); + let toolkit = cuda_toolkit_root(|var| std::env::var(var).ok()); + let header_candidates = cuda_header_candidates(&toolkit, std::env::consts::ARCH); + match header_candidates.iter().find(|path| path.is_file()) { + Some(found) => println!("✓ {}", found.display()), + None => { + println!("✗ not found in the CUDA toolkit at `{}`", toolkit); + eprintln!(" Probed:"); + for candidate in &header_candidates { + eprintln!(" {}", candidate.display()); + } + eprintln!(" Host crates (cuda-bindings) cannot build without cuda.h. Set"); + eprintln!(" CUDA_TOOLKIT_PATH or CUDA_HOME to a CUDA Toolkit install root;"); + eprintln!(" when neither is set, /usr/local/cuda is used."); + ok = false; + } } - // 4. CUDA toolkit + // 5. CUDA toolkit print!("CUDA toolkit (nvcc)... "); match Command::new("nvcc").arg("--version").output() { Ok(output) if output.status.success() => { @@ -878,7 +953,7 @@ pub fn doctor(ctx: &Context) { } } - // 4b. libNVVM + nvJitLink + libdevice (only required when a kernel uses + // 5b. libNVVM + nvJitLink + libdevice (only required when a kernel uses // CUDA libdevice math, e.g. sin/cos/exp/pow). All three ship with the // CUDA Toolkit; checking them here surfaces missing or split packagings // before a runtime failure inside `cuda_host::ltoir::load_kernel_module`. @@ -924,7 +999,7 @@ pub fn doctor(ctx: &Context) { } } - // 5. llc (LLVM static compiler for PTX) + // 6. llc (LLVM static compiler for PTX) // // cuda-oxide requires LLVM 21+: earlier releases reject modern TMA / // tcgen05 / WGMMA intrinsic signatures. Probe in the same order as the @@ -1023,7 +1098,7 @@ pub fn doctor(ctx: &Context) { } } - // 6. clang / libclang resource dir (host `cuda-bindings` / bindgen) + // 7. clang / libclang resource dir (host `cuda-bindings` / bindgen) // // The host `cuda-bindings` crate's build.rs runs bindgen, which loads // libclang at runtime to parse `wrapper.h`. That parse pulls in @@ -1064,7 +1139,34 @@ pub fn doctor(ctx: &Context) { } } - // 7. cuda-gdb (optional) + // 8. NVIDIA driver / GPU. Informational, not fatal: only `cargo oxide + // run` (kernel execution) needs a driver. Cross-compiling and GPU-less + // CI boxes are supported workflows (`build`/`pipeline` work fine), and + // the examples-compile CI job is exactly that. + print!("NVIDIA driver / GPU... "); + match query_gpu_name_and_compute_cap() { + Some((name, (major, minor))) => { + println!("✓ {} (compute capability {}.{})", name, major, minor); + } + None => { + // Some containers mount the kernel driver without shipping + // nvidia-smi; /proc distinguishes "driver loaded, tool broken" + // from "no driver at all". + if Path::new("/proc/driver/nvidia/version").exists() { + println!("- driver loaded, but nvidia-smi is missing or not reporting a GPU"); + eprintln!(" A kernel-mode NVIDIA driver is present (/proc/driver/nvidia/"); + eprintln!(" version), but `nvidia-smi` did not report a usable GPU."); + eprintln!(" `cargo oxide run` may still work; arch auto-detection will fall"); + eprintln!(" back to the backend default (override with --arch=)."); + } else { + println!("- no NVIDIA driver detected"); + eprintln!(" Only `cargo oxide run` (kernel execution) needs the driver;"); + eprintln!(" `cargo oxide build` and `pipeline` work without one."); + } + } + } + + // 9. cuda-gdb (optional) print!("cuda-gdb (optional)... "); match Command::new("cuda-gdb").arg("--version").output() { Ok(output) if output.status.success() => { @@ -1089,6 +1191,43 @@ pub fn doctor(ctx: &Context) { } } +/// CUDA toolkit install root for doctor's `cuda.h` probe: the first set +/// variable among `CUDA_TOOLKIT_PATH`, `CUDA_HOME`, else `/usr/local/cuda`. +/// +/// Kept in lockstep BY HAND with `crates/cuda-bindings/build.rs` +/// (`cuda_toolkit_dir` / `find_cuda_include_dir` / `toolkit_target_dir`): +/// doctor cannot import that probe because build.rs logic is not a library, +/// and cuda-bindings is the NVIDIA-proprietary crate cargo-oxide must not +/// depend on. If the build.rs discovery changes, mirror it here. +fn cuda_toolkit_root(mut get_env: impl FnMut(&str) -> Option) -> String { + ["CUDA_TOOLKIT_PATH", "CUDA_HOME"] + .iter() + .find_map(|var| get_env(var)) + .unwrap_or_else(|| "/usr/local/cuda".to_string()) +} + +/// Candidate `cuda.h` paths under `toolkit`, in probe order: the standard +/// `include/` layout first, then the redistributable `targets/

/include` +/// layout. CUDA names the target dirs after the GPU platform, not the Rust +/// triple: x86_64 hosts use `x86_64-linux`, aarch64 servers use `sbsa-linux`. +/// +/// `arch` is the host CPU architecture; the caller passes +/// `std::env::consts::ARCH` (doctor runs at runtime, so there is no cargo +/// `TARGET` to consult). Injected as a parameter for unit tests. +fn cuda_header_candidates(toolkit: &str, arch: &str) -> Vec { + let base = Path::new(toolkit); + let mut candidates = vec![base.join("include/cuda.h")]; + let target_dir = match arch { + "x86_64" => Some("x86_64-linux"), + "aarch64" => Some("sbsa-linux"), + _ => None, + }; + if let Some(dir) = target_dir { + candidates.push(base.join("targets").join(dir).join("include/cuda.h")); + } + candidates +} + // ============================================================================= // Setup command // ============================================================================= @@ -1437,6 +1576,27 @@ fn parse_compute_cap_field(field: &str) -> Option<(u32, u32)> { Some((major.parse().ok()?, minor.parse().ok()?)) } +/// Query the name and compute capability of the first GPU via `nvidia-smi`, +/// for doctor's driver / GPU report. Same trust rules as +/// [`query_device_compute_cap`]. +fn query_gpu_name_and_compute_cap() -> Option<(String, (u32, u32))> { + let output = Command::new("nvidia-smi") + .args(["--query-gpu=name,compute_cap", "--format=csv,noheader"]) + .output() + .ok() + .filter(|o| o.status.success())?; + parse_gpu_name_and_compute_cap(&String::from_utf8_lossy(&output.stdout)) +} + +/// Parse the first line of `nvidia-smi --query-gpu=name,compute_cap` output +/// into the GPU name and `(major, minor)` pair. Splits on the LAST comma: +/// GPU names may contain commas in principle, `compute_cap` never does. +fn parse_gpu_name_and_compute_cap(stdout: &str) -> Option<(String, (u32, u32))> { + let line = stdout.lines().next()?; + let (name, cap) = line.rsplit_once(',')?; + Some((name.trim().to_string(), parse_compute_cap_field(cap)?)) +} + /// Format a `(major, minor)` compute-capability tuple as the `sm_XX` / /// `sm_XXX[a]` string the codegen backend expects on `CUDA_OXIDE_TARGET`. /// @@ -1971,6 +2131,61 @@ mod tests { assert_eq!(parse_compute_cap("9.0\n12.0\n"), Some((9, 0))); } + #[test] + fn parse_gpu_name_and_compute_cap_splits_on_last_comma() { + assert_eq!( + parse_gpu_name_and_compute_cap("NVIDIA GeForce RTX 5090, 12.0\n"), + Some(("NVIDIA GeForce RTX 5090".to_string(), (12, 0))) + ); + // Failure banner: no comma-separated cc field. + assert_eq!( + parse_gpu_name_and_compute_cap("NVIDIA-SMI has failed.\n"), + None + ); + assert_eq!(parse_gpu_name_and_compute_cap(""), None); + } + + #[test] + fn cuda_toolkit_root_prefers_toolkit_path_then_home_then_default() { + let toolkit_and_home = cuda_toolkit_root(|var| match var { + "CUDA_TOOLKIT_PATH" => Some("/cuda/toolkit".to_string()), + "CUDA_HOME" => Some("/cuda/home".to_string()), + _ => None, + }); + assert_eq!(toolkit_and_home, "/cuda/toolkit"); + + let home_only = + cuda_toolkit_root(|var| (var == "CUDA_HOME").then(|| "/cuda/home".to_string())); + assert_eq!(home_only, "/cuda/home"); + + assert_eq!(cuda_toolkit_root(|_| None), "/usr/local/cuda"); + } + + #[test] + fn cuda_header_candidates_cover_standard_and_redistributable_layouts() { + // Standard install layout first, then the matching targets/ layout. + assert_eq!( + cuda_header_candidates("/usr/local/cuda", "x86_64"), + vec![ + PathBuf::from("/usr/local/cuda/include/cuda.h"), + PathBuf::from("/usr/local/cuda/targets/x86_64-linux/include/cuda.h"), + ] + ); + // aarch64 servers use the sbsa-linux target dir. + assert_eq!( + cuda_header_candidates("/opt/ctk", "aarch64"), + vec![ + PathBuf::from("/opt/ctk/include/cuda.h"), + PathBuf::from("/opt/ctk/targets/sbsa-linux/include/cuda.h"), + ] + ); + // Unknown host arch: only the standard layout is probed. + assert_eq!( + cuda_header_candidates("/opt/ctk", "riscv64"), + vec![PathBuf::from("/opt/ctk/include/cuda.h")] + ); + } + #[test] fn parse_compute_cap_rejects_failure_banners_and_garbage() { // nvidia-smi prints failure text to STDOUT, not stderr. diff --git a/crates/cargo-oxide/src/main.rs b/crates/cargo-oxide/src/main.rs index 229f60f70f..907df62a68 100644 --- a/crates/cargo-oxide/src/main.rs +++ b/crates/cargo-oxide/src/main.rs @@ -210,7 +210,9 @@ fn main() { commands::scaffold_new(&name, async_mode); } Commands::Doctor => { - let ctx = commands::resolve_context(); + // Side-effect-free resolver: doctor must never build the backend + // (or clone anything) before it can diagnose the environment. + let ctx = commands::resolve_doctor_context(); commands::doctor(&ctx); } Commands::Setup => { diff --git a/cuda-oxide-book/getting-started/installation.md b/cuda-oxide-book/getting-started/installation.md index 34bef593e1..cad48ca2f0 100644 --- a/cuda-oxide-book/getting-started/installation.md +++ b/cuda-oxide-book/getting-started/installation.md @@ -229,7 +229,16 @@ Run the built-in diagnostics check: cargo oxide doctor ``` -`cargo oxide doctor` validates your Rust toolchain, CUDA toolkit (including libNVVM / nvJitLink / libdevice for kernels that use math intrinsics), LLVM installation, and codegen backend in one shot. +`cargo oxide doctor` validates your Rust toolchain, CUDA headers (`cuda.h`), +CUDA toolkit (including libNVVM / nvJitLink / libdevice for kernels that use +math intrinsics), LLVM installation, NVIDIA driver / GPU, and codegen backend +in one shot. + +`cargo-oxide` itself does not need the CUDA toolkit or an NVIDIA driver to +build and run, so `doctor` works even on a machine where nothing is installed +yet and diagnoses exactly what is missing (for example missing `cuda.h`, or +no driver). The driver / GPU check is informational: only `cargo oxide run` +needs a GPU, while `build` and `pipeline` work without one. Then build and run an example end-to-end: From 5df7daa24a3610a0f8b3d760a1fb379ea79e88e1 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Fri, 12 Jun 2026 15:46:03 +0530 Subject: [PATCH 4/6] ci: drop the CUDA toolkit from the cargo-oxide unit-test job cargo-oxide no longer pulls cuda-bindings, so its test job needs neither the toolkit install (build-time bindgen) nor the libcuda.so.1 stub shadow. The job now runs on a stock GPU-less, toolkit-less runner, which makes it the exact environment of issue #87; a follow-up step runs the built binary's doctor there and asserts the cuda.h diagnosis is reached (doctor exits 1 by design on that runner, so only the output is asserted). --- .github/workflows/unit-tests.yml | 17 +++++++++++++++-- 1 file changed, 15 insertions(+), 2 deletions(-) diff --git a/.github/workflows/unit-tests.yml b/.github/workflows/unit-tests.yml index dccffdf9a1..86c535ee6d 100644 --- a/.github/workflows/unit-tests.yml +++ b/.github/workflows/unit-tests.yml @@ -52,9 +52,10 @@ jobs: # pure-Rust kernel-launch helpers + tcgen05 tiling tables - package: cuda-host needs_cuda: true - # cargo subcommand wrapper; transitively pulls cuda-bindings + # cargo subcommand wrapper; pure Rust, no CUDA toolkit required. + # Running it on a runner with NO toolkit installed is the + # regression gate for issue #87. - package: cargo-oxide - needs_cuda: true # async DeviceOperation pipeline; transitively pulls cuda-bindings - package: cuda-async needs_cuda: true @@ -106,3 +107,15 @@ jobs: fi cargo test -p ${{ matrix.package }} ${{ matrix.test_args || '--all-targets' }} + + # End-to-end proof for issue #87: on this runner there is no CUDA + # toolkit and no driver, yet doctor must build, load, and reach its + # checks (it exits 1 here because nvcc/cuda.h are missing, which is + # exactly the diagnosis it exists to print). + - name: cargo oxide doctor reaches its checks without the CUDA toolkit + if: matrix.package == 'cargo-oxide' + run: | + set -uo pipefail + out="$(cargo run -p cargo-oxide -- doctor 2>&1)" || true + printf '%s\n' "${out}" + printf '%s\n' "${out}" | grep -q "CUDA headers (cuda.h)" From 76a33c6a1d7b5a9a30131846515fc006551bf3f3 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Fri, 12 Jun 2026 15:46:30 +0530 Subject: [PATCH 5/6] ci: drop the libcuda stub shadow from examples-compile The shadow existed solely because the cargo-oxide binary carried a DT_NEEDED libcuda.so.1 (its run-mode GPU-arch probe used the driver API). cargo-oxide now probes via nvidia-smi and links no CUDA at all, so the binary loads on a GPU-less runner unaided. Example binaries are still linked against the toolkit's lib64/stubs at link time and are never executed in --compile-only mode, so nothing else on the runner needs the soname shadow. Kept as its own commit so it reverts cleanly if some future job step starts exec-ing a cuda-linked binary. --- .github/workflows/examples-compile.yml | 13 ------------- 1 file changed, 13 deletions(-) diff --git a/.github/workflows/examples-compile.yml b/.github/workflows/examples-compile.yml index b894224236..056dde8f55 100644 --- a/.github/workflows/examples-compile.yml +++ b/.github/workflows/examples-compile.yml @@ -80,19 +80,6 @@ jobs: export LIBRARY_PATH="${sysroot}/lib:${LIBRARY_PATH:-}" export LD_LIBRARY_PATH="${sysroot}/lib:${LD_LIBRARY_PATH:-}" - # cargo-oxide itself links libcuda (it can query the GPU arch in - # `run` mode), so its DT_NEEDED carries libcuda.so.1. The dev - # sub-packages only ship the link-time stub named libcuda.so; - # shadow it under the soname so the binary loads on a GPU-less - # runner. Compile-only never makes a real driver call. Same - # trick as unit-tests.yml. - if [ -f "${CUDA_TOOLKIT_PATH}/lib64/stubs/libcuda.so" ]; then - stub_shadow="${RUNNER_TEMP}/libcuda-stub" - mkdir -p "${stub_shadow}" - ln -sf "${CUDA_TOOLKIT_PATH}/lib64/stubs/libcuda.so" "${stub_shadow}/libcuda.so.1" - export LD_LIBRARY_PATH="${stub_shadow}:${LD_LIBRARY_PATH}" - fi - scripts/smoketest.sh --compile-only --no-color - name: Upload failure logs From 9735b13cee2d9c6096da893ab3c5d3dbc0706927 Mon Sep 17 00:00:00 2001 From: nihalpasham Date: Fri, 12 Jun 2026 15:49:01 +0530 Subject: [PATCH 6/6] docs(book): update the doctor checks table in the debugging chapter Doctor gained CUDA headers (cuda.h) and driver / GPU checks, the backend-.so check became informational, and doctor no longer builds the backend before running. Reflect all of that in the checks table and add a short note on which checks are informational. --- .../error-handling-and-debugging.md | 27 ++++++++++++------- 1 file changed, 18 insertions(+), 9 deletions(-) diff --git a/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md b/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md index 4930c239e3..15a048c5a8 100644 --- a/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md +++ b/cuda-oxide-book/gpu-programming/error-handling-and-debugging.md @@ -177,15 +177,17 @@ cargo oxide doctor Doctor checks: -| Check | What it verifies | -|:----------------|:------------------------------------------------------------------------------| -| Rust toolchain | Nightly compiler with required components | -| CUDA toolkit | `nvcc` found and version compatible | -| libNVVM | `libnvvm.so` (CUDA Toolkit) loadable -- needed for libdevice math kernels | -| nvJitLink | `libnvJitLink.so` (CUDA Toolkit) loadable -- needed for libdevice math kernels| -| libdevice | `libdevice.10.bc` discoverable -- needed for libdevice math kernels | -| LLVM | `llc` (21+) available for PTX generation | -| Codegen backend | `librustc_codegen_cuda.so` found (run `cargo oxide setup` to build it) | +| Check | What it verifies | +|:----------------|:-----------------------------------------------| +| Rust toolchain | Nightly compiler with required components | +| Codegen backend | `librustc_codegen_cuda.so` built | +| CUDA headers | `cuda.h` present under the toolkit root | +| CUDA toolkit | `nvcc` found and version compatible | +| libNVVM | `libnvvm.so` loadable (libdevice math kernels) | +| nvJitLink | `libnvJitLink.so` loadable (same) | +| libdevice | `libdevice.10.bc` discoverable (same) | +| LLVM | `llc` (21+) available for PTX generation | +| Driver / GPU | `nvidia-smi` reports a GPU and its compute cap | The libNVVM / nvJitLink / libdevice checks fire only when a kernel calls CUDA libdevice math (`sin`, `cos`, `exp`, `pow`, `sqrt`, ...). If your @@ -193,6 +195,13 @@ kernel is pure arithmetic, those three failing is harmless. They all ship with the CUDA Toolkit -- no separate download. If any check fails, doctor prints the standard install location for that component. +Doctor itself needs neither the CUDA toolkit nor a driver, and it never +builds anything first, so it works on a machine where nothing is installed +yet. Two checks are informational rather than fatal: the codegen backend (a +missing `.so` just means "run `cargo oxide setup`"; `run`/`build` build it +on demand anyway) and the driver / GPU check (only `cargo oxide run` needs +a GPU; `build` and `pipeline` work without one). + ## `cargo oxide pipeline` -- inspecting the compilation When a kernel produces wrong results but no errors, inspect the compilation