Rust GPU Offload: Portable, Safe, and Fast - Paper Analysis

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.

What Rust Offload Solves

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 Three Interfaces

The work provides three offload interfaces with increasing control:

Interface A: Automatic Data Movement

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.

Interface B: Vendor Library Interoperability

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.

Interface C: Explicit Memory Staging

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.

Safe GPU Kernels Without Raw Pointers

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.

The Two-Pass Compilation Pipeline

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.

Performance: Near Parity, With Caveats

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):

• MI250X: Rust is 32% faster to 43% slower than RAJA variants
• H100: Rust is 11% faster to 46% slower than base CUDA

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.

The Optimization Problem

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:

1. 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.

2. Loop-invariant code motion (LICM): Hoist data transfers out of loops. Handle loop-unrolled patterns (same kernel, different scalar offsets) by canceling intermediate transfers.

3. 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.

Critical Assessment

Strengths:

• The Region/PartitioningStrategy design is the most principled approach to safe GPU kernels to date. It avoids raw pointers without compromising generality, and it is extensible by users.
• Two-pass compilation is the right call for a systems language that embeds target-specific code. The alternative of single-pass compilation makes cfg semantics incorrect for either host or device.
• Building on LLVM Offload (rather than inventing a new runtime) is pragmatic. It inherits multi-device fat binary support and ongoing upstream optimization work.
• The three-interface design (convenient / library / explicit) covers the spectrum from prototyping to production.

Weaknesses and Open Questions:

• The 400x gap between Interface A and C is a usability trap. Most users will start with Interface A, get terrible performance, and conclude Rust GPU programming is slow. The compiler optimizations to close this gap are prototyped but not benchmarked.
• Memory transfer overhead (46 ms vs 16 ms for RAJA on H100) is unexplained and significant. This suggests the async transfer prototyping is not yet production-ready.
• Register pressure (33 vs 28) from bounds-checking will compound with larger, register-bound kernels. For compute-bound HPC workloads, this could be a real penalty.
• Cross-pass metadata serialization adds complexity. If a crate uses complex generic hierarchies, the metadata round-trip between passes could introduce subtle bugs that are hard to diagnose.
• Shared memory is exposed as raw pointers with no safety guarantees. The paper acknowledges this explicitly but it means tile-based algorithms (reductions, matrix multiply) remain unsafe.
• No multi-GPU support yet. The paper mentions it as future work.

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.

References

• Paper: https://arxiv.org/abs/2608.13759
• Authors: Manuel S. Drehwald, Marcelo Dominguez, Kevin Sala, Alan Aspuru-Guzik, Johannes Doerfert
• Affiliations: University of Toronto, LLNL, Universidad Rey Juan Carlos