COOLJAPAN
← All posts

OxiCUDA 0.5.1 Released — oxicuda-nvrtc Completes the Zero-SDK Runtime-JIT Story

OxiCUDA 0.5.1 adds oxicuda-nvrtc, a pure-Rust runtime loader for NVIDIA's NVRTC CUDA-C-to-PTX JIT compiler, completing the zero-SDK-dependency story alongside oxicuda-driver. Graceful degradation, process-wide caching, and direct PTX handoff to oxicuda-driver. 38,646 tests, ~1.30M SLoC, 74 crates.

release oxicuda cuda gpu pure-rust nvrtc jit ptx

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:

OxiCUDA 0.5.1 ends all of that.

Technical Deep Dive

  1. Foundation — oxicuda-driver, oxicuda-memory, oxicuda-launch, oxicuda-runtime, and now oxicuda-nvrtc. Dynamic driver loading via libloading, an RAII DeviceBuffer<T>, and the type-safe launch! macro already sat underneath everything else; oxicuda-nvrtc joins this layer as the runtime-JIT half of the same “zero SDK dependency” promise oxicuda-driver established for the driver API.
  2. PTX codegen & autotuner — oxicuda-ptx, oxicuda-autotune. ElementwiseTemplate and ReductionTemplate hand-generate PTX text directly; oxicuda-nvrtc offers 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.
  3. 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.
  4. Scientific computing and the wider domain layer. oxicuda-fft, oxicuda-sparse, oxicuda-solver, and oxicuda-rand cover 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

Tips

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

↑ Back to all posts