baracuda

About the name. Yes, we know — it's spelled barracuda (two Rs). That name was taken on crates.io, so we dropped one R and kept swimming.

A unified Rust ML-op facade over the NVIDIA CUDA ecosystem.

What baracuda is

baracuda is a Rust workspace that exposes every primitive an ML framework expects — the union of PyTorch (torch.* + nn.functional) and JAX (jax.lax.* + jax.numpy.*) — through a single Plan-based crate surface called baracuda-kernels. Internally each plan dispatches to:

1. The appropriate NVIDIA-library wrapper crate (cuBLAS, cuDNN, cuFFT, cuSOLVER, cuRAND, cuSPARSE, cuTENSOR, NPP, CV-CUDA, CUTLASS) when one already covers the op well, or
2. A bespoke hand-rolled .cu kernel shipped in baracuda-kernels-sys when no NVIDIA library covers the op (or covers it poorly at the shapes that matter for modern transformer / vision / GNN workloads).

Callers import one crate (baracuda-kernels) and reach for one API style. The dispatch decision — which is observable through Plan::sku() for telemetry — is otherwise invisible. Switching from a CUTLASS-backed SKU to a bespoke-backed SKU is a layout flag, not an import change.

baracuda is for downstream Rust ML / inference / training frameworks that need access to the full CUDA stack without re-vendoring it themselves. The workspace also ships idiomatic stand-alone wrappers for every CUDA library under crates/baracuda-<lib> if you want to skip the kernel facade and talk to one library directly.

Status

In active development — alpha.67. 2280+ GPU tests passing, zero failures across the 6 critical test crates on an RTX 4070 (sm_89; the baracuda-kernels suite alone is 2180/0 across 513 test binaries). alpha.67 (Phase 74, Fuel-ask) ships the plain dense FP GEMM family: 12 cuBLAS-backed flat C symbols (baracuda_kernels_gemm_dense_{f32, f64, f16, bf16}_*) with RRR / RCR / CRR layouts, flexible leading dims, and strided-batch folded into the base signature — plus the DenseGemmPlan<T> typed plan, the ReduceToPlan<T, N> broadcast-reverse reduction facade, UnaryKind::Step, and gelu flavor-disambiguation docs. This closes the last non-baracuda CUDA surface in Fuel (its own cuBLAS MatMul wrapper). alpha.66 (Fuel-ask) exposed per-device VRAM queries on the Driver-API Device — vram_free() / vram_total() / vram_info() wrap cuMemGetInfo_v2 so downstream optimizers can read runtime memory pressure without round-tripping through a failed allocation. Phase 73 follow-up (alpha.65) lands a 17-33× decode speedup via a focused FlashDecodingPlan (split-K, seq_q=1), 4× win at GQA shapes via the new num_kv_heads descriptor field, the long-awaited FlashSdpaPlan GQA-broadcast routing fix, and a SDPA gap closure that makes baracuda's standard MHA shape 50% faster than PyTorch by making fa2 a default feature. The ConcatPlan (KV-cache decode) and reduce_axis (small-shape rows) kernels also got significant rewrites — 13× and 2.6-15.4× respectively. Phase 63 (alpha.63, Fuel-ask) closes the FlashAttention saved-tensor wiring gap: NEW baracuda_kernels_fa2_sdpa_lse_size(batch, num_heads, seq_q) -> usize dense LSE pre-allocation helper (sibling of the existing varlen _varlen_lse_size); load-bearing "LSE saved-tensor contract" docs on the FW + BW trailblazers naming the exact FW→saved-LSE→BW handoff; new docs guide at docs/guides/fa2-saved-tensor-contract.md showing the wiring pattern downstream autograd frameworks should use. Test investment: 8 new tests — 3 host-only lse_size helper sanity tests + 5 GPU FW→BW roundtrip tests (f16/bf16 × d128 causal/noncausal + d64 causal) + 4 BW feature-surface tests backfilling the gaps left by fa2_backward_smoke.rs (BW with sliding window, softcap, ALiBi, and all-features composed). The existing softmax_lse output on baracuda_kernels_fa2_sdpa_<dt>_run (v1) and ..._run_v2 (Phase 59a) — turns out Fuel didn't need a new "v3 with lse" variant; the FW has been writing LSE since alpha.56, just lacked the size helper + saved-tensor wiring documentation. Phase 62 (alpha.62, Fuel-ask) lifts the in-place op contract from contig-only (Phase 61) to strided by shipping 11 new affine in-place FFI symbols (4 contig int dtype backfill — i32/i64/u8/i8 matching forward affine; 7 strided variants across the full forward-strided dtype matrix — f32/f64/i32/i64/u8/bf16/f16) and documenting the stride-equality precondition for same-pointer aliasing on the unary / binary / ternary strided trailblazers as a stable public contract. NEW baracuda_kernels_types::strides_equal host helper for callers to validate the precondition before dispatch. Zero new bespoke kernels for the elementwise unary/binary/ternary families — their existing strided launchers are aliasing-safe under the contract. Test investment: 14 host-only unit tests + 17 GPU smoke tests for the new FFI surface + 7 aliasing-contract proof tests across contig + strided trailblazers (also backfills the Phase 61 contig contract that shipped without test coverage). Phase 61 (alpha.61) completed the alpha.55 baseline with bf16/f16 in-place affine + the original contig same-pointer contract doc. Phase 60 lifted FA2 FW to the full Candle-fork-extended 9-head_dim set ({32, 64, 96, 128, 160, 192, 224, 256, 512}) via 12 new vendored .cu files + 32 FW smoke tests. Phase 59a + 59b added the full FA2 v2.8.3 surface (FW + BW + varlen across head_dims 32-256, GQA, ALiBi, sliding window, softcap) plus 48 new smoke tests, closing Fuel's FA2-retirement requirements. Phase 59c (consolidation pass, alpha.59) fixed a pre-existing parallel-test race in the bespoke flash kernel's SMEM-carveout call surfaced by Phase 59a's 5-head_dim fanout, plus updated flash_sdpa_backward_smoke to force the bespoke backend on f16/bf16 (Phase 59b made FA2 the new default BW backend, breaking source-compat for the existing bespoke BW smoke tests). Phase 42-44 add three opt-in backends (FA2, mHC.cu, ozIMMU); none are on the default build path. Phase 44b internalized the ozIMMU sources (clean-fork; cutf submodule retired; Linux + Windows both build). Phase 49 adds the baracuda-optim sibling crate (Adam / LAMB / SGD via vendored Apex multi_tensor_apply) — gated behind the optim feature so inference-only consumers don't pay the FFI surface cost. Phase 55 adds the baracuda-transformer-engine sibling crate (NVIDIA TransformerEngine FP8 cast + delayed-scaling recipe, Apache-2.0) — gated behind the tensor_engine feature. On Ada (sm_89) the FP8 wins are bandwidth-saving only (KV cache, weights); the recipe machinery is forward-compatible with Hopper / Blackwell where the MMA throughput win also materializes.

Phase coverage (see ARCHITECTURE.md for the phase matrix):

API stability is not promised before beta.0. Breaking changes ship in each alpha bump and are documented in the workspace CHANGELOG.md.

Quick start

Add the kernel facade and the driver crate:

[dependencies]

baracuda-kernels = { version = "0.0.1-alpha.64", features = ["sm89", "cudnn"] }

baracuda-driver  = "0.0.1-alpha.64"

A representative example — single-axis numerically stable softmax over a device-resident tensor:

use baracuda_driver::{Context, Device, DeviceBuffer, Stream};
use baracuda_kernels::{
    PlanPreference, SoftmaxArgs, SoftmaxDescriptor, SoftmaxKind, SoftmaxPlan,
    TensorMut, TensorRef, Workspace,
};

fn main() -> Result<(), Box<dyn std::error::Error>> {
    // 1. Standard CUDA bring-up via baracuda-driver.
    let ctx = Context::new(&Device::get(0)?)?;
    let stream = Stream::new(&ctx)?;

    // 2. Allocate device input + output buffers (rank-2: rows × cols).
    let rows = 32i32;
    let cols = 1024i32;
    let n_elems = (rows * cols) as usize;
    let dev_x: DeviceBuffer<f32> = DeviceBuffer::zeros(&ctx, n_elems)?;
    let mut dev_y: DeviceBuffer<f32> = DeviceBuffer::zeros(&ctx, n_elems)?;

    // 3. Build the descriptor — pure shape + dtype + op-kind, no handles.
    let desc = SoftmaxDescriptor::<2> {
        kind: SoftmaxKind::Softmax,
        input_shape: [rows, cols],
        softmax_axis: 1,
        element: <f32 as baracuda_kernels::KernelDtype>::KIND,
    };

    // 4. Plan selection — picks a kernel SKU (bespoke softmax kernel here).
    let plan = SoftmaxPlan::<f32, 2>::select(&stream, &desc, PlanPreference::default())?;

    // 5. Args carry the per-call tensor handles + strides.
    let args = SoftmaxArgs {
        x: TensorRef { data: dev_x.as_slice(), shape: [rows, cols], stride: [cols as i64, 1] },
        y: TensorMut { data: dev_y.as_slice_mut(), shape: [rows, cols], stride: [cols as i64, 1] },
    };

    // 6. Launch. Workspace::None for plans that need no scratch.
    plan.run(&stream, Workspace::None, args)?;
    stream.synchronize()?;
    Ok(())
}

The same select → run shape applies to every op. GEMM, attention, conv2d, FFT, scatter — the descriptor / args fields differ per family but the lifecycle is identical. See the crates/baracuda-kernels README for the int8-GEMM variant of the same example.

Workspace layout

The user-facing crates a typical caller will reach for:

baracuda-kernels             # the unified Plan-based ML op facade
baracuda-kernels-types       # shared type vocabulary (Element, TensorRef, KernelSku, ...)
baracuda-kernels-sys         # raw FFI to bespoke .cu kernels
baracuda-kernels-bench       # criterion harness for sm_89 perf sweeps (not published)
baracuda-cutlass             # safe wrapper for CUTLASS GEMM (float, int8 RCR, batched, grouped)
baracuda-driver              # safe wrapper for the CUDA Driver API
baracuda-runtime             # safe wrapper for the CUDA Runtime API

The per-library wrappers used internally by the facade (you can also use them stand-alone):

baracuda-cublas{,-sys}       # cuBLAS + cuBLASLt + cuBLASXt
baracuda-cudnn{,-sys}        # cuDNN classic + Graph API
baracuda-cudf{,-sys}         # cuDF (RAPIDS dataframe; Linux-only)
baracuda-cufft{,-sys}        # cuFFT
baracuda-cusolver{,-sys}     # cuSOLVER dense + sparse + Rf + Mg
baracuda-cusparse{,-sys}     # cuSPARSE
baracuda-curand{,-sys}       # cuRAND
baracuda-cutensor{,-sys}     # cuTENSOR
baracuda-cutlass{,-sys}      # CUTLASS GEMM kernel templates
baracuda-cutlass-kernels-sys # CUTLASS kernel-only compile target
baracuda-cuvs{,-sys}         # RAPIDS cuVS GPU vector search (Phase 71)
baracuda-cvcuda{,-sys}       # CV-CUDA image processing
baracuda-flashinfer{,-sys}   # FlashInfer paged-KV + cascade + sampling (Phase 46/66)
baracuda-npp{,-sys}          # NPP
baracuda-nccl{,-sys}         # NCCL
baracuda-nvcomp{,-sys}       # nvCOMP
baracuda-nvimagecodec{,-sys} # nvImageCodec (Phase 70 — supersedes nvJPEG)
baracuda-nvjpeg{,-sys}       # nvJPEG (kept for back-compat)
baracuda-nvshmem{,-sys}      # NVSHMEM symmetric-heap RDMA (Phase 69)
baracuda-ozimmu{,-sys}       # ozIMMU Ozaki-scheme DGEMM (Phase 44)
baracuda-transformer-engine{,-sys}  # TransformerEngine FP8 (Phase 55)

And the supporting low-level crates (FFI, build infrastructure, profiling):

baracuda-cuda-sys            # Driver + Runtime FFI
baracuda-nvrtc{,-sys}        # runtime CUDA C++ → PTX
baracuda-nvjitlink{,-sys}    # CUDA 12+ JIT linker
baracuda-cupti{,-sys}        # profiling APIs
baracuda-nvml{,-sys}         # device monitoring
baracuda-cufile{,-sys}       # GPUDirect Storage (Linux-only)
baracuda-tensorrt{,-sys}     # TensorRT inference runtime (Phase 68 — vtable-dispatch C++ shim)
baracuda-forge              # build-time .cu → PTX compiler driver
baracuda-build              # build.rs helpers
baracuda-core                # loader + Error plumbing
baracuda-types{,-derive}     # pure-data types: Half, BFloat16, Complex, DeviceRepr

Vendor-track training / optimization crates (opt-in via feature gates on baracuda-kernels):

baracuda-megatron            # Megatron-LM TP primitives (Phase 57; composition, no kernels)
baracuda-optim               # NVIDIA Apex multi-tensor optimizers (Phase 49)

The full umbrella crate (baracuda) re-exports everything behind cargo features — convenient when you want everything; overkill when you don't.

Hardware support

baracuda targets Ampere and newer by design. Pre-Ampere GPUs lack the tensor-core instructions and async-copy primitives the bespoke kernels are written against (mma.sync.m16n8k*, cp.async, ldmatrix), and we have no desire to ship a slower SIMT fallback for hardware that's eight years old.

The default sm80 build runs forward-compatibly on Ada and Hopper through JIT-compiled PTX; turn on sm89 to pick up the FP8 and Flash-Attention sibling plans tuned for Ada's larger register file.

Cargo features

The kernel facade exposes a broad opt-in feature set. Architecture and library-integration features are off by default — pick what your deployment needs. Important: unless a feature is enabled, its plans are not built, not linked, and not present in the public API. Always check the feature gate when scanning for an op family.

Architecture targets

NVIDIA library integrations

Vendored kernel families

Sibling crates also have their own feature gates:

• baracuda-cuvs — cuvs feature in baracuda-kernels-sys ecosystem pulls in RAPIDS cuVS vector-search (Phase 71, Linux-only). Not yet re-exported by baracuda-kernels itself.
• baracuda-tensorrt — shim feature builds the vtable-dispatch C++ shim required to call TensorRT (Phase 68 — no flat C ABI upstream).
• baracuda-flashinfer — same flashinfer feature as above.

Notes:

• cudnn is off by default because cuDNN is a separate NVIDIA download not bundled with the stock CUDA toolkit installer. Enabling it without cuDNN installed produces a linker error on cudnn.lib / libcudnn.so — see the Building section for auto-discovery paths.
• Most vendored features add 30s–5 min to first build (template-heavy CUDA). Subsequent builds incremental.
• Features that pull in optional sibling crates (optim, tensor_engine, megatron_tp, nvshmem, ring_attention) only compile the sibling when the feature is enabled — inference-only consumers don't pay the surface cost.

Building

Requirements:

• CUDA Toolkit ≥ 12.0 with nvcc on PATH. baracuda is tested on 12.x and 13.x.
• cuDNN 9.x (only if you enable the cudnn feature) — separate NVIDIA download, not bundled with the toolkit.
• A working Rust toolchain ≥ 1.85 (workspace MSRV pinned in rust-toolchain.toml).
• Windows users: lld-link.exe somewhere on PATH. The CUDA nvcc invocation links through it; the install location is typically C:\Program Files\LLVM\bin. Install the LLVM Windows package and add that directory to PATH if cargo build complains about lld-link.exe not being found.

A typical full build with all GPU-side features (CUDA toolkit + cuDNN present):

cargo build -p baracuda-kernels --features sm89,cudnn --release

Or, to verify the public API surface compiles without the full kernel build (fast — type-check only):

cargo check -p baracuda-kernels --features sm89,cudnn

The baracuda-kernels-sys build script auto-discovers cuDNN at the following paths in order: CUDNN_PATH / CUDNN_ROOT / CUDNN_HOME env vars, then C:\Program Files\NVIDIA\CUDNN\v<X.Y>\ on Windows, then the CUDA toolkit's own lib/ directory (pre-cuDNN-9 layout), then the standard Linux distro paths under /usr/lib/.

Troubleshooting

Windows: Git-for-Windows fake link.exe shadowing the MSVC linker

Git-for-Windows ships a GNU coreutils binary named link.exe at C:\Program Files\Git\usr\bin\link.exe — its job is to create a hard link, not to link object files. If that directory appears on PATH ahead of the MSVC linker (or LLVM's lld-link.exe), cargo build invokes the coreutils binary instead of the real linker and fails with a cryptic error (it doesn't understand /OUT: and friends).

baracuda's baracuda-kernels-sys and baracuda-cutlass-sys build scripts probe PATH on Windows and emit a cargo:warning if they detect this shadowing. Fix: re-order PATH so the MSVC linker (typically reached via the Visual Studio "x64 Native Tools Command Prompt") or LLVM's lld-link.exe (C:\Program Files\LLVM\bin\) appears before C:\Program Files\Git\usr\bin\. Building from the VS x64 Native Tools prompt is the most reliable option; alternatively, install LLVM and put its bin directory ahead of Git's on the user/system PATH.

Testing

baracuda's GPU integration tests are gated behind #[ignore] so a host-only cargo test doesn't try to launch a kernel on a machine without an NVIDIA driver. To run them you need a working GPU plus the --ignored flag:

# Host-only tests (compile + reference logic; no GPU access):

cargo test -p baracuda-kernels --lib


# Full GPU integration sweep — RTX 30xx / 40xx / 50xx required:

cargo test -p baracuda-kernels --release -- --ignored


# Verify the workspace-level API surface compiles (no GPU needed):

cargo check -p baracuda-kernels --features sm89,cudnn

The full regression on an RTX 4070 covers 324 binary targets at ~1630 tests passing. Individual op-family suites take 30–90 seconds; the full sweep is 25–40 minutes.

Benchmarks

The baracuda-kernels-bench crate is a criterion-based harness with CUDA-event-timed throughput sweeps across GEMM, Flash Attention, and Conv2d at LLM-typical and ResNet-typical shapes. It is not published to crates.io (it depends on a working GPU).

cargo bench -p baracuda-kernels-bench --features sm89,cudnn

The full sweep takes ~30 minutes on an RTX 4070. Scope to a single family with --bench gemm / --bench flash_attention / --bench conv2d. See crates/baracuda-kernels-bench/BENCH-sm89.md for the baseline table format and methodology.

Project documentation

• ARCHITECTURE.md — layered design, Plan-Descriptor-Args pattern, KernelSku taxonomy, dispatcher design, workspace contract, sibling-plan pattern, vendoring convention, phase roadmap.
• OP-MATRIX.md — full op × dtype × backend coverage matrix (planned).
• LESSONS.md — postmortems, ABI footguns, performance traps (planned).
• Per-crate README.md files under crates/<name>/.

License

Dual-licensed under MIT or Apache-2.0. Pick whichever fits your project. Contributions accepted under the same terms.

NVIDIA's CUDA libraries (libcuda, libcudart, libcublas, libcudnn, …) are not redistributed by this project. You obtain them from NVIDIA separately — either through the CUDA Toolkit installer or through each library's dedicated download page. baracuda's loader opens whatever the host driver / toolkit has installed.

Vendor attribution

A small number of bespoke kernels in baracuda-kernels-sys are vendored from upstream open-source projects (huggingface/candle's CUDA kernel set via fuel-cuda-kernels; llama.cpp's ggml-cuda GGUF block-format quantization + MMVQ; guoqingbao/attention.rs's fused MoE expert kernels). Each adapted source carries an SPDX-FileCopyrightText: + SPDX-License-Identifier: header; the consolidated provenance is in crates/baracuda-kernels-sys/LICENSE-thirdparty.md.

FlashAttention v2 (Tri Dao, BSD-3-Clause, pinned at v2.8.3 / 060c9188beec3a8b62b33a3bfa6d5d2d44975fab) is vendored at crates/baracuda-kernels-sys/vendor/flash-attention/ with verbatim LICENSE + AUTHORS files and full vendor / scope notes in VENDOR.md. Gated behind the fa2 cargo feature on baracuda-kernels-sys and baracuda-kernels; exposed through a backend-choice path on FlashSdpaPlan (Phase 42).

mHC.cu (Andre Slavescu, MIT, pinned at a426939c2dbc11c443db041bcff12b65d1b6482a) — unofficial CUDA implementation of DeepSeek-AI's Manifold-Constrained Hyper-Connections paper — is vendored at crates/baracuda-kernels-sys/vendor/mhc/ with the verbatim upstream LICENSE, an AUTHORS file, and full vendor / scope notes in VENDOR.md. Gated behind the mhc cargo feature on baracuda-kernels-sys and baracuda-kernels; exposed through the new HyperConnectionPlan (Phase 43, Tier 1: static-H FW, bf16 weights / f32 activations).

FlashInfer (NVIDIA

• FlashInfer community, Apache-2.0 with full patent grant, pinned at v0.6.12 / eee0d75f91f64c520bfaed07e39a850ea4ddde23) — a curated ~12 kLOC subset of the FlashInfer header tree is vendored at crates/baracuda-kernels-sys/vendor/flashinfer/ with verbatim upstream LICENSE + NOTICE and full vendor / scope / patch notes in VENDOR.md. Gated behind the flashinfer cargo feature on baracuda-kernels-sys and baracuda-kernels; exposes three NEW plan families — paged-KV decode + append (BatchPagedDecodePlan + PagedKvAppendPlan for vLLM-style serving), sort-free sampling (TopKTopPSamplingPlan — combined TopK/TopP/MinP via a single-kernel rejection sampler), and cascade attention LSE merge (CascadeAttentionPlan for prefix-cache sharing). Surgical cherry-pick (not a wholesale wrap) — Hopper / Blackwell / NVSHMEM / Mamba / MLA / POD paths intentionally skipped to keep the build cost contained. Phase 46.

The baracuda-forge build-time kernel-compiler crate is a vendored fork of cudaforge by Guoqing Bao — see crates/baracuda-forge/NOTICE for the upstream commit hash.

The baracuda-cutlass safe wrapper for NVIDIA CUTLASS — plan-based GEMM and grouped-GEMM with caller-supplied workspace, MoE-friendly variable-M-per-group dispatch — was specified by the Fuel ML library team. See crates/baracuda-cutlass/NOTICE for the design lineage.