cudarc: a safe Rust surface over the CUDA toolkit, split into sys, result and safe

A safe layer over an API that cannot be made entirely safe

The CUDA toolkit is a C API. Calling it from Rust by hand means writing FFI bindings, managing device memory lifetimes, and getting buffer sizes right, which is exactly the class of code where a mistake is a memory fault rather than a compile error. cudarc is a Rust crate that provides safe wrappers over the CUDA libraries, and the project describes itself as a minimal and safe api over the cuda toolkit.

The project states its design goals plainly, and the first one is the interesting one: as safe as possible, with the honest caveat that there will still be a lot of unsafe due to ffi and async. That sentence tells you what you are buying. You are not getting a proof of safety, you are getting a layer that keeps the unsafe parts in the places where the CUDA API actually forces them, and you are expected to stay on the safe side of that line.

The second goal is ergonomics, and the third is the one that shapes the whole API: allow mixing of high level `safe` apis with low level `sys` apis. That is why the crate is not a single flat surface.

The audience is Rust developers who need GPU compute without writing their own bindings, and the library covers the driver API, NVRTC, cuRAND, cuBLAS, cuBLASLt, NCCL, cuDNN, cuSPARSE, cuSOLVER, cuFILE, CUPTI, NVTX and cuFFT.

Every wrapper comes in sys, result and safe

The layering is the central idea. Each library wrapper exposes three levels, and by default the safe API is what you get:

use cudarc::driver::{safe, result, sys};
use cudarc::nvrtc::{safe, result, sys};
use cudarc::cublas::{safe, result, sys};
use cudarc::cublaslt::{safe, result, sys};
use cudarc::curand::{safe, result, sys};
use cudarc::nccl::{safe, result, sys};

The three levels are defined in the documentation. `sys` is the raw ffi apis generated with bindgen. `result` is a very small wrapper around sys to return `Result` from each function. `safe` is a wrapper around result and sys to provide safe abstractions.

The consequence is that a failure in a CUDA call is a `Result` rather than an error code you have to remember to check, and that a resource such as a context or a stream is a type with a lifetime rather than a raw handle you might drop twice. The cost of the three layer scheme is verbosity at the boundary, since a project that needs one specific cuBLAS call not covered by the safe layer has to drop down to `result` or `sys` and accept the safety trade.

The project recommends sticking with the safe APIs. The bindings themselves are produced by a `bindings_generator/` directory in the repository, and `build.rs` drives the build.

The CUDA version is a build time feature, not a runtime choice

Because the CUDA libraries are versioned and the FFI signatures change between releases, cudarc makes the toolkit version a compile time decision. The supported ranges are listed: CUDA 11.4 to 11.8, 12.0 to 12.9, and 13.0 to 13.4. cuDNN is supported at 8.9.7, 9.10.2 and 9.21.1, and NCCL at 2.22 to 2.30.

The selection is done with a Cargo feature flag, for example `-F cuda-13010` for a specific version. There is also `-F cuda-version-from-build-system`, which gets the CUDA toolkit version at build time using `nvcc`, and a companion `-F fallback-latest` that controls the behaviour if that detection fails. The documentation is explicit that the fallback is not enabled by default, and that without it the build script will panic. That is a defensible choice, since silently compiling against the newest bindings when the local toolkit is older produces a program that compiles and then fails at load time.

The manifest shows this as one feature per version, from `cuda-11040` up to `cuda-13040`, and the `docs.rs` metadata enables a specific set of them so the published documentation is built against a known combination.

Dynamic loading by default, and a support matrix that is not uniform

The second build time decision is how the libraries are linked. By default the crate uses `-F dynamic-loading`, which will not require any libraries to be present at build time. That is the sensible default for a crate that most people will add to a project without a CUDA installation on the build machine. `-F dynamic-linking` and `-F static-linking` are available for other cases.

The support table is where the constraints show. The CUDA driver supports dynamic load and dynamic link but not static link. NVRTC, cuRAND, cuBLAS, cuBLASLt, NCCL, cuDNN, cuSPARSE, cuFILE and CUPTI support all three, while cuFFT is marked as supporting only dynamic loading. cuSOLVER and NVTX are marked as not applicable for static link.

So the linking mode you want is not always available. A build that needs cuFFT statically linked is not something this crate can do, and that is a packaging decision you have to make before you start rather than discover in CI.

The default feature set in the manifest makes the same choice: `std`, `cublas`, `cublaslt`, `curand`, `driver`, `runtime`, `nvrtc` and `fallback-dynamic-loading` are on by default, while `cusparse`, `cusolver`, `cufile`, `nvtx`, `cupti` and `cufft` are opt in.

Context, stream, copy, and the first kernel launch

The getting started section is short enough to reproduce. Creating a context and moving data to the device looks like this:

// Get a stream for GPU 0
let ctx = cudarc::driver::CudaContext::new(0)?;
let stream = ctx.default_stream();

// copy a rust slice to the device
let inp = stream.clone_htod(&[1.0f32; 100])?;

// or allocate directly
let mut out = stream.alloc_zeros::<f32>(100)?;

Kernels are launched through a builder that takes the arguments one at a time:

let mut builder = stream.launch_builder(&sin_kernel);
builder.arg(&mut out);
builder.arg(&inp);
builder.arg(&100usize);
unsafe { builder.launch(LaunchConfig::for_num_elems(100)) }?;

And the data comes back to the host with:

let out_host: Vec<f32> = stream.clone_dtoh(&out)?;
assert_eq!(out_host, [1.0; 100].map(f32::sin));

The `alloc_zeros` and `clone_htod` calls show where the safety lives. The slice is borrowed for the duration of the copy, the device buffer owns its allocation, and the returned `Vec` is a plain Rust value. The launch, by contrast, is marked `unsafe`, which is the honest boundary: the launch configuration and the argument buffer lengths are not checked by the compiler.

Compiling a kernel at runtime with NVRTC

cudarc also exposes NVRTC, so a kernel can be written as source text in your program, compiled to PTX at runtime, and loaded into the context without a separate build step:

let ptx = cudarc::nvrtc::compile_ptx("
extern \"C\" __global__ void sin_kernel(float *out, const float *inp, const size_t numel) {
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < numel) {
        out[i] = sin(inp[i]);
    }
}")?;

// Dynamically load it into the device
let module = ctx.load_module(ptx)?;
let sin_kernel = module.load_function("sin_kernel")?;

The pipeline is source string, PTX, module, function, and then the same builder used for any other kernel. That is what makes the crate useful for a library that wants to ship a kernel without requiring the user to have nvcc, and it is also the reason NVRTC appears in the default feature set alongside the driver.

The trade is that the kernel source becomes a string in your binary rather than a compiled object, so the compiler cannot check it, and errors in it surface at runtime. The `examples/` directory in the repository shows the same idea as files, with `sin.cu` and `sin.ptx` checked in alongside `nvrtc-compile.rs`, and numbered examples covering allocation, copies, kernel launches, streams, device representation, threading, CUPTI, constant memory, function attributes, context configuration and multi-GPU copies.

Where the safety stops, and where the support stops

Three limits deserve to be stated without hedging.

First, the crate does not remove unsafe code. The design goals say there will still be a lot of unsafe due to ffi and async, and the kernel launch in the example is an `unsafe` block. A team that requires a codebase free of unsafe blocks will not find that here.

Second, the linking matrix is incomplete in the ways listed earlier, and cuFFT in particular supports dynamic loading only. A pipeline that needs a statically linked cuFFT has to look elsewhere.

Third, the version support is a moving target that cudarc cannot lead. The crate ships bindings for specific toolkit releases, and each new CUDA release means a new set of features and a new cudarc release to enable them. The release notes show the pattern: v0.19.8 was titled for cuda 13.3, and v0.19.10 for Cuda 13.4 and quality of life. If you pin a CUDA version in your container, you are pinning it in your Cargo features too, and the two have to be updated together.

One more detail is worth flagging for a no-std user. The manifest lists `no-std` among the crate categories and the `std` feature is in the default set, so the standard library is opt out, not opt in.

Dual licensed, and a 0.x version line that still moves quickly

The licence is dual, and the project states the reason: it is dual-licensed to be compatible with the Rust project. The text is the Apache License, Version 2.0 or the MIT license, at your option, and the repository carries `LICENSE-APACHE` and `LICENSE-MIT` at the root. The manifest agrees, with `license = "MIT OR Apache-2.0"`. For a crate that many Rust projects depend on, that is the expected choice and removes any licensing question from an evaluation.

On maintenance, the crate is pre-1.0 and moving. v0.19.10 was published on 2026-09-24 and titled for Cuda 13.4 and quality of life, v0.19.9 on 2026-08-11 with bugfixes, and v0.19.8 on 2026-06-19 for cuda 13.3. The last push to the repository was on 2026-09-24, the same day as the newest tag.

A 0.x version means the API can still change between releases, so pinning an exact version in a Cargo.lock is the sensible approach, and reading the release notes when you bump it is not optional. The project also maintains a Discord server and the documentation on docs.rs, with the crate published on crates.io.

Editorial conclusion

cudarc fits a Rust project that has to call GPU libraries and wants the compiler to check lifetimes and buffer handling instead of hand written FFI, and it fits teams that need to pick a CUDA version at build time rather than at runtime. It does not fit a project that needs cuFFT through dynamic linking or static linking, since the support table marks those combinations unavailable, and it does not remove unsafe code from a kernel launch, which the documentation leaves as an unsafe call. Before you pin a version, note that the crate tracks specific toolkit releases, from CUDA 11.4 through 13.4, so a new toolkit version means waiting for a new cudarc release, and the latest one, v0.19.10, added CUDA 13.4 support.

Frequently asked questions

Official sources

1. chelsea0x3b/cudarc on GitHub
2. Issues
3. License: Apache-2.0
4. README
5. Releases