Rust GPU Offload: Portable, Safe, and Fast - Paper Analysis
rustgpucudahipllvmparallel-computingmemory-safetyopenmphpc
Loading…
In August 2026, Manuel Drehwald (University of Toronto / LLNL), Marcelo Dominguez (Universidad Rey Juan Carlos), Kevin Sala (LLNL), Alan Aspuru-Guzik (University of Toronto), and Johannes Doerfert presented Rust Offload: a compiler extension to rustc that lets you write GPU kernels in Rust and offload them to NVIDIA or AMD GPUs through the LLVM Offload infrastructure. The paper is arXiv 2608.13759.
This is not another Rust GPU experiment. It is a path toward vendor-agnostic GPU computing in a memory-safe language, built directly into the compiler.
The Rust GPU ecosystem has been fragmented. rust-gpu targets SPIR-V and emulates pointers, making it unsuitable for HPC. rust-cuda is NVIDIA-only and requires raw pointers for all mutable arguments, abandoning safety. cuda-oxide is NVIDIA-led and uses single-pass compilation that conflates host and device target semantics. None of these achieve both portability and safety simultaneously.
Rust Offload does. It compiles a single Rust codebase to both NVIDIA (PTX) and AMD (AMDGCN) targets, and it provides a novel abstraction that allows safe GPU kernels using standard Rust slices instead of raw pointers.
The work provides three offload interfaces with increasing control:
offload!(kernel_1, &input, &mut output);
Rust's reference types determine data directionality automatically: &T becomes a read-only device transfer, &mut T becomes bidirectional. The compiler generates the OpenMP target data-mapping clauses from the type system. Scalars up to 64 bits are passed by value.
This is the easiest to use but has a critical performance pitfall: if the host reads output between two kernel calls, the compiler must insert a device-to-host and host-to-device synchronization, because it cannot prove the transfer is redundant. In the paper's benchmarks, naive usage of this interface is over 400x slower than the explicit interface.
core::intrinsics::offload_args::<_, _, ()>(
rocblas_sgemv_wrapper,
(&A, &x, &mut y)
);
Wraps cuBLAS, rocBLAS, and other vendor libraries as offload operations. The compiler can reason about data transfers across the call boundary using the same type-based infrastructure. This lets programmers incrementally replace CPU hotspots with GPU library calls.
let out_gpu = core::intrinsics::preload_mut(&mut output);
offload!(kernel_1, &input, &out_gpu);
offload!(kernel_2, &input, &out_gpu);
drop(out_gpu); // explicit sync back to host
PreloadMut<'a, T> borrows the host value (enforced by the borrow checker) and stages it on the device. The GPU handle holds a raw pointer internally (not a reference), because after preloading the live data lives at a different address on the device. A PhantomData<&'a mut T> maintains the lifetime/aliasing connection without creating UB. Dropping the handle triggers the device-to-host transfer.
This eliminates the hidden synchronization problem of Interface A while remaining safe: the borrow checker prevents host access to the value while the device holds it.
The core challenge: Rust's aliasing rules say a mutable reference is unique, but GPU kernels share mutable references across threads. rust-cuda solves this by requiring raw pointers for all mutable arguments.
Rust Offload instead introduces Region<T, PartitioningStrategy>, which partitions data into disjoint subregions per thread:
let mut reg = Region::<_, Linear1D>::new(&mut x, ());
PartitioningStrategy is an unsafe trait (implementations must guarantee disjointness), but the resulting kernel code is safe Rust. Users write standard slice operations; the partitioning strategy computes thread indices internally and returns non-overlapping mutable references.
This design is inspired by cuda-oxide's DisjointSlice, but differs in three ways: (1) it is a standalone crate users can extend with custom partitioning schemes, (2) it supports returning chunks of data, not just scalar elements, and (3) it is simpler -- no special ThreadIndex type needed.
This is the most consequential architectural decision in the paper.
The compiler runs twice. Pass 1: rustc targets the GPU architecture (amdgcn-amd-amdhsa or nvptx64). It compiles kernel code and produces a device binary. Pass 2: rustc targets the host (x86_64-unknown-linux-gnu). It compiles the application code and embeds the device binary into the host LLVM IR, lowering offload intrinsics to libomptarget runtime calls.
Why two passes instead of one (like cuda-oxide)? Because Rust is a systems language with #[cfg(target_arch = "x86_64")] guards, inline assembly, and target-specific intrinsics (AVX-512, NEON). If you compile with the CPU target for the frontend, cfg evaluation selects CPU-specific code paths even for what will become GPU kernels. If you compile with the GPU target, host-side optimizations are lost. Two passes let each target evaluate cfg independently.
The cost: monomorphization tracking breaks between passes. The host's main is compiled in pass 2, but kernel instantiations are needed in pass 1. The solution is cross-pass metadata exportation: pass 2 serializes kernel DefIds and concrete type substitutions, and pass 1 imports them to seed monomorphization.
This is the same approach used by OpenMP offload, HIP, and CUDA. The toolchain is built on LLVM 23.1.0-rc1. The paper argues that single-pass designs either (a) conflate target semantics or (b) require carrying unevaluated cfg attributes through the entire compiler, which would pessimize compile times for non-GPU users.
The paper ports RAJAPerf (the benchmark suite for RAJA, a C++ portable parallelism framework) to pure Rust and compares against RAJA-CUDA and RAJA-HIP on MI250X, H100, and RTX A2000.
Kernel times: Near-parity. Rust kernels perform similarly to RAJA kernels across all benchmarks, except FIR and LTIMES (small micro-benchmarks sensitive to unrolling decisions).
Total runtime (excluding memory transfers):
Memory transfers: Rust moves less data overall (423 MB vs 468 MB on H100) but takes longer (46 ms vs 16 ms). The paper attributes this to differences in memory kinds and async transfer handling -- the async transfer prototype is not yet production-ready.
Register usage: Rust averages 33 registers vs RAJA's 28 on RTX 2070. Likely caused by bounds-checking on explicit thread/block indexing.
Fast-math: Rust's algebraic floats provide a 2x speedup on FIR (small loop that vectorizes at width 4) and ~20% on several other kernels. Notably, Rust does not support C++-style fast-math because the nnan/ninf flags can trigger UB in safe Rust code. The algebraic float approach exposes most optimization opportunities without those assumptions.
Interface A is the natural entry point but can be 400x slower than Interface C. The paper prototypes three compiler optimizations to close this gap:
Async data transfers: Split each offload! into separate H2D, launch, and D2H calls. Start H2D as early as possible, start D2H immediately after kernel completion, block only when the data is needed.
Loop-invariant code motion (LICM): Hoist data transfers out of loops. Handle loop-unrolled patterns (same kernel, different scalar offsets) by canceling intermediate transfers.
Clippy lint: Warn when a preload call is placed unnecessarily late, wasting time between last CPU usage and H2D transfer start.
The paper claims these optimizations can make Interface A match Interface C performance on RAJAPerf, but does not present completed benchmark data for this.
Strengths:
Weaknesses and Open Questions:
What This Means for the Rust HPC Ecosystem:
Rust Offload is the first proposal that does not force you to choose between safety and portability. cuda-oxide gives you safety but only on NVIDIA. rust-gpu gives you portability but with pointer emulation. rust-cuda gives you neither in a principled way. Rust Offload gives you both, with competitive kernel performance.
The question is whether the LLVM Offload dependency is a feature or a liability. It provides immediate vendor support and ongoing optimization work, but it also means the project's fate is tied to LLVM's offload infrastructure, which has historically been under-maintained and under-documented. The authors note they are engaging with upstream maintainers, which is encouraging.
The paper does not address compilation time. Two-pass compilation means every build invokes rustc twice per GPU target. For large codebases, this could be a significant developer experience issue. The metadata serialization is described as "negligible" but has not been benchmarked.