◉ COOLJAPAN
← All posts
July 7, 2026 • 10 min read

OxiCUDA 0.4.1 Released — Concurrency Races, Cross-Backend Correctness, and Real CUDA Graph Capture

OxiCUDA 0.4.1 extends on-device GPU validation into BLAS, DNN, and the driver/memory/launch stack, catching concurrency races that single-threaded testing can't. Plus the first correctness pass across five non-CUDA backends and real driver-backed CUDA Graph capture. 38,612 tests, ~1.28M SLoC, 73 crates.

release oxicuda cuda gpu pure-rust concurrency vulkan webgpu cuda-graphs

A kernel that passes validation running alone on the GPU can still be wrong the moment two streams touch it at the same time.

Today we released OxiCUDA 0.4.1 — a release that pushes the on-device validation sweep started in 0.4.0 into new territory. It’s no longer just “does this kernel compute the right numbers on real silicon” — it’s “does this still hold up under real concurrent execution, and does it hold up on hardware that was never CUDA’s to begin with.”

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 more, across 73 crates and ~1.28M lines of safe Rust. libcuda.so / nvcuda.dll is loaded dynamically at runtime via libloading; nothing about building OxiCUDA requires the CUDA SDK or a C/C++ compiler. The result compiles to a single static binary and runs everywhere the driver is present — from Turing through Blackwell, and, through OxiCUDA’s own Vulkan, Metal, WebGPU, ROCm, and Level Zero backends, on GPUs that were never NVIDIA’s to begin with.

Why OxiCUDA 0.4.1 is a game changer

NVIDIA’s official stack still asks you to accept, on faith:

OxiCUDA 0.4.1 ends all of that.

Technical Deep Dive: what 0.4.1 actually validates

  1. Foundation correctness — driver, runtime, memory, launch. The sweep reached the layer everything else stands on: oxicuda-driver’s context/stream/event/module lifetime (a new live-context registry stops a child handle like a Stream or Module from double-freeing the driver when it outlives the Context it was created from), oxicuda-runtime’s error propagation (device_synchronize was silently swallowing real driver errors instead of returning them), and oxicuda-memory’s pool/peer/managed-memory paths (cross-device allocation, peer-access context restoration, and uninitialized PinnedBuffer/UnifiedBuffer/MappedBuffer memory now zeroed before it ever reaches safe code).
  2. PTX codegen hardening — oxicuda-ptx. New validators (RegisterAllocator::validate_named(), validate_branch_labels(), tensor-core operand-consistency checks) catch illegal or colliding register declarations and mismatched branch targets before a single instruction is emitted, layered on top of ISA-level fixes to SM90+ warp/TMA primitives (mbarrier, elect.sync, fence.proxy, cp.async.bulk.tensor) and tensor-core MMA/WGMMA encoding.
  3. Scientific and deep-learning kernels — oxicuda-blas, oxicuda-dnn, oxicuda-sparse, oxicuda-solver, oxicuda-autotune. Each crate now carries its own gpu-tests on-device validation suite, JIT-compiling the crate’s actual production PTX via Module::from_ptx and diffing device output against an independent CPU oracle — the pattern that caught oxicuda-solver’s batched LU/Cholesky kernels being literal ret; stub bodies, now rewritten as real barrier-staged Doolittle LU and Cholesky factorizations.
  4. Multi-vendor backends — oxicuda-vulkan, oxicuda-webgpu, oxicuda-metal, oxicuda-rocm, oxicuda-levelzero. Each backend’s SPIR-V, WGSL, MSL, or HIP generation now gets checked by that backend’s own real compiler or validator, and every backend’s GEMM entry point now rejects transpose/leading-dimension configurations it can’t actually honor instead of silently computing the wrong answer.

Getting Started

Add OxiCUDA and opt into the subsystems you need:

cargo add oxicuda --features blas

Default features are driver, memory, and launch. Every subsystem is its own flag (ptx, autotune, blas, dnn, fft, sparse, solver, rand, primitives, vulkan, metal, webgpu, rocm, level-zero, and full for everything) — plus gpu-tests if you have real GPU hardware and want to run the same on-device validation this release relies on.

A complete GEMM, end to end:

use oxicuda::prelude::*;

fn main() -> Result<(), oxicuda::Error> {
    // Initialize driver and select GPU device
    let device = Device::get(0)?;
    let ctx = Context::new(device)?;
    let stream = Stream::new(&ctx)?;

    // Allocate device memory
    let mut d_a = DeviceBuffer::<f32>::zeroed(1024)?;
    let mut d_b = DeviceBuffer::<f32>::zeroed(1024)?;
    let mut d_c = DeviceBuffer::<f32>::zeroed(1024)?;

    // Copy host data to device
    d_a.copy_from_host(&host_a)?;
    d_b.copy_from_host(&host_b)?;

    // Launch a GEMM: C = alpha * A @ B + beta * C
    let handle = BlasHandle::new(&stream)?;
    handle.gemm(
        Transpose::None, Transpose::None,
        m, n, k,
        1.0f32,            // alpha
        &d_a, lda,
        &d_b, ldb,
        0.0f32,            // beta
        &mut d_c, ldc,
    )?;

    stream.synchronize()?;

    // Copy result back to host
    let mut result = vec![0.0f32; m * n];
    d_c.copy_to_host(&mut result)?;
    Ok(())
}

On a machine with an NVIDIA GPU, run the same validation this release depends on:

cargo test --features gpu-tests

No nvcc invocation, no .cu files, no linker flags chasing a CUDA install. cargo build, and you ship.

What’s New in 0.4.1

Tips

This is the foundation

OxiCUDA is the GPU layer beneath the rest of the COOLJAPAN ecosystem, and 0.4.1 is already the version every sibling project depends on. SciRS2 pulls in oxicuda-driver/-memory/-fft/-ptx/-launch/-blas/-solver/-sparse/-dnn; ToRSh routes its CudaBackend through oxicuda-backend/-driver/-launch/-ptx; sklears carries oxicuda-backend/-memory/-blas/-solver/-manifold/-dnn/-driver/-ptx/-primitives; OxiONNX executes graphs through oxicuda-driver/-memory/-blas/-dnn/-ptx/-launch. OxiFFT, OxiML, OxiProj, QuantRS2, and TensorLogic — which builds its own tensorlogic-oxicuda-* sparse/solver/rng bridge crates on top — round out the list. Every one of them is already pinned to oxicuda* = "0.4.1": the concurrency, cross-backend, and CUDA Graph fixes in this release are live under all of them today.

Repository: https://github.com/cool-japan/oxicuda

Star the repo if you believe GPU correctness should be proven under real concurrency and on real hardware, not assumed from a CPU stand-in or a single-threaded test run once. Every star tells us to keep building.

Pure Rust GPU computing is here — fast, safe, sovereign, and now proven under the conditions it actually has to survive.

— KitaSan at COOLJAPAN OÜ
July 7, 2026

↑ Back to all posts