◉ COOLJAPAN
← All posts
August 12, 2026 • 8 min read

OxiCUDA 0.5.4 Released — WarpVec/WarpMask Bring Simd-Style Vectors to CUDA Warps

OxiCUDA 0.5.4 adds WarpVec/WarpMask, a SIMD-flavored warp-vector layer for oxicuda-ptx's builder DSL, six new IR instructions for warp shuffle/vote, and a foreign-compiler PTX interop test proving the driver/launch stack hosts upstream rustc-generated PTX too. 38,689 tests, ~1.30M SLoC, 74 crates.

release oxicuda cuda gpu pure-rust simd ptx warp-programming rustc-interop

A CUDA warp is already a 32-lane vector unit. OxiCUDA 0.5.4 is the first release to let you program it like one.

Today we released OxiCUDA 0.5.4 — a feature release adding WarpVec/WarpMask, a Simd/Mask-style expression layer over oxicuda-ptx’s builder DSL that treats a CUDA warp’s 32 lanes as a first-class vector value, plus the low-level warp shuffle/vote IR instructions it’s built on, and a new foreign-compiler PTX interop test proving the driver/launch stack correctly hosts rustc-generated PTX modules, not just oxicuda-ptx’s own.

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 74 crates and ~1.30M lines of safe Rust. libcuda.so/nvcuda.dll is 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.4 is a game changer

Warp-level programming is where CUDA’s ergonomics fall apart fastest:

OxiCUDA 0.5.4 ends all of that.

Technical Deep Dive

  1. IR layer — oxicuda-ptx. Six new instructions back the warp-vector layer: Shfl (shfl.sync, all four source-lane modes via the new ShflMode enum), Vote (vote.sync, all/any/uni/ballot via VoteMode), Mov (typed register/immediate moves with hex-exact float immediates), PackB64x2/UnpackB64x2 (routing 64-bit values through the 32-bit-only shuffle datapath), and Not (incl. not.pred) — plus MulMode::Rn for explicit mul.rn.f32/.f64. All six are fully wired into the validator’s def/use analysis, dead-code elimination, register-pressure tracking, instruction scheduling, arch-legality checks (sm_70 floor), and the interactive TUI explorer’s category/latency model.
  2. Emission layer — builder::warp_ops. The low-level primitives under WarpVec/WarpMask: lane_id, warp_index_x, mov_typed, shfl_sync/shfl_sync_with_valid (CUDA-exact segmented c-operand encoding), vote_{all,any,uni,ballot}, and not_pred.
  3. Validation layer — oxicuda-primitives. 11 on-device gpu-tests, validated on an RTX A4000 (sm_86), exercise WarpVec/WarpMask against independent CPU oracles: relu-dot, the redux.sync fast path, s32/f64/segmented reductions, min/max, both scan directions, votes/ballot/count/segmented-any, and every shuffle mode.
  4. Interop layer — oxicuda-launch’s rustc_ptx_interop. JIT-compiles and launches PTX fixtures produced by upstream nightly rustc’s own NVPTX backend (nvptx64-nvidia-cuda) — not oxicuda-ptx’s generator — directly through the driver stack: a scalar SIMT saxpy and a core::simd (portable-SIMD) relu-dot kernel, both checked for exact numerical agreement against a CPU oracle. Worth being precise here: the portable-SIMD fixture documents empirically that upstream rustc scalarizes Simd within a single thread — no shfl.sync is emitted by rustc itself. The lane-to-warp mapping that WarpVec provides above is oxicuda-ptx’s own codegen layer, layered on top; it isn’t something stable Rust’s core::simd gives you across a CUDA warp today.

Getting Started

cargo add oxicuda-ptx

A complete warp-level relu-dot-product kernel, generated entirely from Rust — no PTX written by hand:

use oxicuda_ptx::prelude::*;

fn main() {
    let ptx = KernelBuilder::new("relu_dot_warp")
        .target(SmVersion::Sm86)
        .param("out", PtxType::U64)
        .param("x", PtxType::U64)
        .param("y", PtxType::U64)
        .body(|b| {
            let x_ptr = b.load_param_u64("x");
            let y_ptr = b.load_param_u64("y");
            let lane = b.lane_id();
            let x_addr = b.f32_elem_addr(x_ptr, lane.clone());
            let y_addr = b.f32_elem_addr(y_ptr, lane.clone());
            let x_val = b.load_global_f32(x_addr);
            let y_val = b.load_global_f32(y_addr);
            let x = WarpVec::from_register(x_val).expect("f32 vec");
            let y = WarpVec::from_register(y_val).expect("f32 vec");

            // Elementwise multiply, ReLU, then a butterfly all-reduce sum
            // across all 32 lanes -- redux.sync on sm_80+, shfl.sync below it.
            let dot = x
                .mul(b, &y).expect("mul")
                .relu(b).expect("relu")
                .reduce_sum(b).expect("reduce");

            // Lane 0 stores the warp's aggregate.
            let out_ptr = b.load_param_u64("out");
            let one = b.mov_imm_u32(1);
            b.if_lt_u32(lane, one, |b| {
                b.store_global_f32(out_ptr.clone(), dot.register().clone());
            });
            b.ret();
        })
        .build()
        .expect("PTX generation failed");

    assert!(ptx.contains("shfl.sync.bfly.b32"));
    println!("Generated {} bytes of PTX using a butterfly warp-shuffle reduction", ptx.len());
}

This is the same kernel shape validated on real hardware. To run the equivalent on your own GPU:

cargo test --features gpu-tests -p oxicuda-primitives

What’s New in 0.5.4

Tips

This is the foundation

OxiCUDA is the GPU layer beneath the rest of the COOLJAPAN ecosystem. Its 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. Every kernel-generating crate built on oxicuda-ptx can reach for WarpVec/WarpMask starting today.

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

Star the repo if you think warp-level parallelism deserves a Simd-style API instead of hand-rolled shfl.sync calls. Every star tells us to keep building.

The era of hand-written warp-shuffle boilerplate is over. Pure Rust GPU computing is here — and as of 0.5.4, your CUDA warps get a Simd-style vector API to match.

— KitaSan at COOLJAPAN OÜ August 12, 2026

↑ Back to all posts