A GPU kernel that ptxas refuses to assemble fails loudly. A GPU kernel that assembles fine and quietly computes the wrong answer is far more dangerous — 0.5.0 found and fixed one of each.
Today we released OxiCUDA 0.5.0 — a correctness release that fixes F64-precision PTX code generation across elementwise and reduction kernels, closes a mixed-precision GEMM bug that was silently corrupting results instead of erroring, and eliminates the last 42 .expect() call sites anywhere in the workspace’s library code.
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.30M 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.
Why OxiCUDA 0.5.0 is a game changer
NVIDIA’s official stack still asks you to accept, on faith:
- Raw CUDA C++ template metaprogramming for every precision variant of a kernel, with no compiler check that a given instantiation is even valid for its target precision
- cuBLAS/cuDNN FFI bindings that hand you a
void*foralpha/betaand trust you to have encoded it at the right precision — get it wrong and you silently get zero or garbage back, not an error - Double-precision code paths that vendor test suites rarely exercise, because most ML workloads never leave F16/BF16/F32 — so F64 bugs can live for releases at a time
.unwrap()/.expect()scattered through a hand-rolled numeric stack, each one a landmine waiting for the one input shape nobody tested
OxiCUDA 0.5.0 ends all of that.
- F64 elementwise kernels stop failing at the assembler. Every elementwise generator — add/sub/mul/div, relu/sigmoid/gelu/silu/tanh, sqrt/rsqrt/exp/log/ceil/floor, pow, min/max, compare, fused-scale-add, fill — hardcoded a value-register name that OxiCUDA’s own auto-declaration logic always sized as 32-bit, so an F64 kernel tried to push 64-bit values through a 32-bit register bank and
ptxasrejected it outright ("Arguments mismatch for instruction 'ld'"). A new precision-awareElementwiseTemplate::vreg_prefix()now picks the correct register class (fdfor F64,fhfor F16/BF16,ffor F32/TF32) — the F32 output is byte-for-byte unchanged. - Transcendental ops now refuse invalid precision instead of silently emitting it.
exp,log,sigmoid,gelu,silu,tanh,softplus, andpowall lower to PTX special-function-unit instructions (ex2.approx/lg2.approx) that only exist at.f32— asking for F64 used to generate a moduleptxaswould reject later, at load time.ElementwiseTemplate::validate_precision()now catches the mismatch at generation time with a descriptive error;rsqrt/sqrt/hard_*/leaky_reluare unaffected since PTX defines those at every precision. - F64 reductions get a real code path. Block-level sum/mean/L2-norm reductions hardcoded an F32 register bank and a 32-bit
shfl.sync.down.b32warp-shuffle for the final stage — physically incapable of moving a 64-bit value in one step. F64 now reduces entirely through a shared-memory tree down to stride 1; the F32 path (tree down to one warp, then warp-shuffle) is unchanged. - Mixed-precision GEMM stops silently zeroing itself.
alpha/betawere bit-encoded at the input precision (F16/BF16/FP8 E4M3/E5M2) instead of the accumulator precision the epiloguealpha * acc + beta * Cactually runs in — so a1.0_f16alpha got reinterpreted as an F32 near-zero denormal, quietly killing the product on every mixed-precision GEMM call. A newGpuFloat::to_accumulator_bits()conversion fixes the encoding for F16, BF16, and both FP8 formats. - Zero
unwrap()/expect()left in library code. The last 42 production.expect()call sites, spread across 22 crates, are gone — replaced with propagatedResult/Optionusing each crate’s own error type, non-panicking float comparisons viatotal_cmp, or structural rewrites that make the invariant hold by construction instead of asserting it at runtime.
Technical Deep Dive
- Foundation —
oxicuda-driver,oxicuda-memory,oxicuda-launch,oxicuda-runtime. Dynamic driver loading vialibloading, an RAIIDeviceBuffer<T>, and the type-safelaunch!macro sit underneath everything else. None of 0.5.0’s fixes touched this layer — it’s been through two prior on-device validation sweeps (0.4.0, 0.4.1) and held steady. - PTX codegen & autotuner —
oxicuda-ptx,oxicuda-autotune. This is where 0.5.0’s core fixes live.ElementwiseTemplateandReductionTemplategenerate the actual PTX text for every elementwise op and block reduction across F16/BF16/F32/TF32/F64, and this release is the first time the F64 path assembles correctly — while also refusing to generate F64 for the handful of ops PTX itself only defines at F32. - 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. The newGpuFloat::to_accumulator_bits()fix lives here, correcting how GEMM’salpha/betaepilogue scalars are encoded across every mixed-precision path. - 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 — all downstream ofoxicuda-ptxandoxicuda-blas, so they inherit this release’s precision fixes without a single line of their own code changing.
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 exercise this release’s F64 and mixed-precision GEMM fixes yourself.
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’s fixes were checked against:
cargo test --features gpu-tests
What’s New in 0.5.0
- F64 elementwise PTX kernels now assemble. Every elementwise generator (arithmetic, activations, unary ops, compare, fused-scale-add, fill) fixed a hardcoded 32-bit register prefix that made
ptxasreject every F64 kernel outright. - Transcendental ops reject invalid precision up front.
exp/log/sigmoid/gelu/silu/tanh/softplus/pownow fail at generation time with a clear error if asked for non-F32 precision, instead of emitting PTX thatptxaswould reject later. - F64 block-level reductions work. Sum/mean/L2-norm and similar reductions now reduce F64 values through a shared-memory tree instead of a 32-bit warp shuffle that can’t carry a 64-bit value.
- Mixed-precision GEMM no longer silently corrupts results.
alpha/betaare now encoded at the accumulator’s F32 precision via a newGpuFloat::to_accumulator_bits()conversion, fixing F16/BF16/FP8 GEMM calls with non-default scalars. bf16_gemm_errorandlora::merge::scale_adapterhonor their documented contracts. The former returns0.0for degenerate (empty-product) dimensions instead of panicking; the latter returnsGenResult<LoraLinear>instead of panicking on a reconstruction failure.- Zero
unwrap()/expect()remain in library code. The last 42 production.expect()call sites, across 22 crates, are gone — completing the workspace’s zero-unwrap policy. - A broken intra-doc link is fixed, so
cargo docno longer fails under-D warnings. - Test suite grew to 38,622 passing tests (
--all-features; 37,296 with default features), up from 38,612/37,288 at 0.4.1.
Tips
- Re-run any F64 elementwise or reduction kernel that previously failed to build. If you hit a
ptxaserror like"Arguments mismatch for instruction 'ld'"generating F64 kernels on 0.4.1 or earlier, 0.5.0 fixes it outright — just upgrade and rebuild. - Don’t request F64 for
exp/log/sigmoid/gelu/silu/tanh/softplus/pow. These lower to F32-only special-function-unit instructions;oxicuda-ptxnow rejects the combination at generation time. Usersqrt/sqrt/hard_*/leaky_reluif you need F64 — they’re unaffected. - Re-check any mixed-precision GEMM output you shipped from 0.4.1 or earlier. F16/BF16/FP8 GEMM calls with non-default
alpha/betamay have been silently near-zeroed by the old input-precision encoding bug; 0.5.0 fixes this with no API changes on your end. - Trust that library code no longer panics on a bad shape. With the last 42
.expect()sites gone, every documentedResult/Optionin OxiCUDA’s public API means what it says — propagate with?rather than coding around a possible panic. bf16_gemm_errornow returns0.0for degenerate dimensions instead of an internal panic — match on the plain value, not aResult, if you were working around this before.- Run
cargo test --features gpu-testson real hardware to confirm these fixes hold on your exact GPU — the on-device validation harness from 0.4.0/0.4.1 covers the same elementwise, reduction, and GEMM kernels this release touched.
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 73 crates, so a single cargo add oxicuda pulls in whichever subsystem a dependent project needs. The F64 PTX fixes and the GEMM precision fix in 0.5.0 are live under every one of them today, without a single line of their own code changing.
Repository: https://github.com/cool-japan/oxicuda
Star the repo if you believe GPU numeric bugs should fail loud at generation time instead of silently shipping a wrong answer to production. Every star tells us to keep building.
Pure Rust GPU computing is here — and as of 0.5.0, its double-precision path finally works as well as its single-precision one.
— KitaSan at COOLJAPAN OÜ July 14, 2026