Hand-written PTX gets you a long way. Compiling arbitrary CUDA-C at runtime — with no CUDA Toolkit installed on the machine that runs it — gets you the rest of the way.
Today we released OxiCUDA 0.5.1 — adding oxicuda-nvrtc, a pure-Rust runtime loader for NVIDIA’s NVRTC library, so CUDA-C source can be JIT-compiled to PTX and launched without the CUDA Toolkit ever touching the target machine.
No CUDA SDK. No nvcc. No C/C++ toolchain. OxiCUDA is a type-safe, memory-safe, pure-Rust replacement for the entire NVIDIA CUDA Toolkit software stack — cuBLAS, cuDNN, cuFFT, cuSPARSE, cuSOLVER, cuRAND, and now NVRTC — across 74 crates and ~1.30M lines of safe Rust. libcuda.so/nvcuda.dll, and now libnvrtc.so/nvrtc64_*.dll, are both loaded dynamically at runtime via libloading; nothing about building or running OxiCUDA requires the CUDA SDK, headers, or a C/C++ compiler. The result compiles to a single static binary and runs everywhere the driver is present.
Why OxiCUDA 0.5.1 is a game changer
NVIDIA’s official NVRTC story still asks you to accept, on faith:
- Linking against
libnvrtc.so/nvrtc.libat build time, which drags the full CUDA Toolkit onto every machine that builds your binary — just to JIT-compile CUDA-C at runtime - A C API that hands back raw
nvrtcResultcodes and leaves log retrieval, header injection, and C++ name-mangling as manual, pointer-juggling multi-step dances - A missing or mismatched NVRTC install that usually surfaces as a link error or a segfault, not a typed, catchable error
- Every other “pure Rust CUDA” story that still bottoms out in an FFI binding assuming the CUDA Toolkit is present wherever the binary ships
OxiCUDA 0.5.1 ends all of that.
- NVRTC loads at runtime, not link time.
oxicuda-nvrtcdlopenslibnvrtc.so(Linux) ornvrtc64_*.dll(Windows) the first time it’s actually needed — no#[link]attribute, nobuild.rs, no-lnvrtc. The crate compiles on any standard Rust toolchain whether or not the CUDA Toolkit is anywhere in sight. - Graceful degradation is the default, not an afterthought. On a host without NVRTC, nothing panics:
is_available()returnsfalse, and every fallible entry point returns a typedNvrtcError::Unavailablenaming the exact library files it tried. NVRTC becomes an optional accelerator you probe once, not a hard runtime dependency. - Optional capabilities degrade independently. CUBIN retrieval, C++ name expressions, and the supported-architecture query each return
NvrtcError::NotSupportedwhen the underlying symbol is missing from an older NVRTC build, instead of failing the whole library load over one optional entry point. - The resolved function table is cached process-wide. A
OnceLockresolves NVRTC’s symbols exactly once per process, so everycompile_to_ptxcall after the first is effectively free. - Output feeds straight back into the rest of OxiCUDA.
Ptx::as_str()returns a NUL-terminated string built to be handed directly tooxicuda_driver::Module::from_ptx— CUDA-C source goes in, a launchable kernel comes out, without ever leaving the workspace.
Technical Deep Dive
- Foundation —
oxicuda-driver,oxicuda-memory,oxicuda-launch,oxicuda-runtime, and nowoxicuda-nvrtc. Dynamic driver loading vialibloading, an RAIIDeviceBuffer<T>, and the type-safelaunch!macro already sat underneath everything else;oxicuda-nvrtcjoins this layer as the runtime-JIT half of the same “zero SDK dependency” promiseoxicuda-driverestablished for the driver API. - PTX codegen & autotuner —
oxicuda-ptx,oxicuda-autotune.ElementwiseTemplateandReductionTemplatehand-generate PTX text directly;oxicuda-nvrtcoffers the complementary path — compiling arbitrary CUDA-C source to PTX at runtime for the cases where hand-written PTX templates aren’t the right tool. - Linear algebra & deep learning —
oxicuda-blas,oxicuda-dnn. Full BLAS L1/L2/L3 with SIMT/Tensor-Core/Split-K GEMM dispatch, plus convolution, attention, normalization, and quantization for DNN — untouched this release, still carrying 0.5.0’s F64 PTX and mixed-precision GEMM fixes. - Scientific computing and the wider domain layer.
oxicuda-fft,oxicuda-sparse,oxicuda-solver, andoxicuda-randcover the cuFFT/cuSPARSE/cuSOLVER/cuRAND surface; above that sit dozens more domain crates — generative diffusion (oxicuda-gen), graph neural networks (oxicuda-gnn), state-space models (oxicuda-mamba), vision transformers (oxicuda-vision), causal inference (oxicuda-causal), PDE solvers (oxicuda-pde), and more — bringing the workspace to 74 crates total.
Getting Started
Add OxiCUDA and opt into NVRTC:
cargo add oxicuda --features nvrtc
Default features remain driver, memory, and launch; nvrtc sits alongside every other subsystem flag (ptx, autotune, blas, dnn, fft, sparse, solver, rand, primitives, vulkan, metal, webgpu, rocm, level-zero, and full for everything) as an opt-in.
Compile CUDA-C to PTX at runtime and feed it straight into the driver:
use oxicuda::nvrtc::{compile_to_ptx, is_available};
fn main() -> Result<(), oxicuda::nvrtc::NvrtcError> {
if is_available() {
let src = r#"
extern "C" __global__ void saxpy(float a, float* x, float* y, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) y[i] = a * x[i] + y[i];
}
"#;
let ptx = compile_to_ptx(src, "saxpy.cu", &["--gpu-architecture=compute_75"])?;
// `ptx.as_str()` feeds `oxicuda_driver::Module::from_ptx` directly —
// no manual FFI, no intermediate file, no CUDA Toolkit required here.
println!("{}", ptx.as_str());
}
Ok(())
}
For finer control — in-memory headers, C++ name expressions, CUBIN output, or inspecting the compiler log — drive a Program directly instead of the compile_to_ptx shortcut; see the oxicuda-nvrtc README for the full API.
On a machine with an NVIDIA GPU and NVRTC installed:
cargo test --features gpu-tests
What’s New in 0.5.1
- New crate:
oxicuda-nvrtc. A pure-Rust runtime loader for NVIDIA’s NVRTC (CUDA-C → PTX JIT compiler), completing the runtime-JIT half of the zero-SDK-dependency story alongsideoxicuda-driver. - Zero build-time CUDA dependency.
libnvrtcisdlopen’d vialibloading— no#[link], nobuild.rs, no-lnvrtc— and the resolved function table is cached process-wide after first use. - Graceful degradation throughout.
is_available()and a typedNvrtcError::Unavailablehandle a missing NVRTC install;NvrtcError::NotSupportedhandles individually-missing optional entry points (CUBIN retrieval, C++ name expressions, supported-arch queries) on older runtimes. - Direct interop with
oxicuda-driver.Ptx::as_str()returns output built to feedoxicuda_driver::Module::from_ptxwith no intermediate step. - Exposed from the umbrella crate as
oxicuda::nvrtcbehind the newnvrtcfeature flag (off by default). - Test suite grew to 38,646 passing tests (
--all-features; 37,316 with default features), up from 38,622/37,296 at 0.5.0 — 20 new tests, bringing the workspace to 74 crates.
Tips
- Reach for
oxicuda-nvrtcwhen hand-written PTX isn’t the right fit. If your kernel is more naturally expressed as CUDA-C than as a hand-built PTX template, compile it at runtime withcompile_to_ptxinstead of hand-writing PTX text. - Always gate on
is_available()first. Treat NVRTC as an optional accelerator: check availability, fall back to a CPU path or pre-built PTX if it’sfalse, and never assume the target machine has an NVIDIA driver and NVRTC installed. - Distinguish
NvrtcError::UnavailablefromNvrtcError::NotSupported. The former means NVRTC itself couldn’t be loaded; the latter means NVRTC loaded fine but a specific optional entry point (CUBIN, name expressions, arch query) isn’t in this runtime’s version — handle them differently rather than treating both as “NVRTC is broken.” - Feed
Ptx::as_str()straight intooxicuda_driver::Module::from_ptx. There’s no intermediate file or manual FFI step between compiling CUDA-C and launching it on the device. - Enable with
--features nvrtc. It’s off by default like every other OxiCUDA subsystem — add it explicitly alongside whatever else your project needs (blas,dnn,ptx, and so on). - Run
cargo test --features gpu-testson real hardware to exerciseoxicuda-nvrtc’s compilation path on your exact GPU and driver combination.
This is the foundation
OxiCUDA is the GPU layer beneath the rest of the COOLJAPAN ecosystem. Its own architecture diagram puts SciRS2, OxiONNX, TrustformeRS, and ToRSh directly on top of it, with OxiBLAS and OxiFFT rounding out the list of projects that lean on this stack — and the oxicuda crate itself is an umbrella re-export over all 74 crates, so a single cargo add oxicuda pulls in whichever subsystem a dependent project needs. oxicuda-nvrtc is live under every one of them today, opt-in and ready, without a single line of their own code changing.
Repository: https://github.com/cool-japan/oxicuda
Star the repo if you believe runtime CUDA-C compilation should be as dependency-free as everything else in a pure-Rust GPU stack. Every star tells us to keep building.
Pure Rust GPU computing is here — and as of 0.5.1, it can JIT-compile CUDA-C itself, with no toolkit required.
— KitaSan at COOLJAPAN OÜ July 22, 2026