When NVIDIA released CUDA (Compute Unified Device Architecture) in 2006, GPUs transformed from components dedicated to graphics rendering into general-purpose parallel computers. Since then, the means of writing kernels (functions that carry out the core computation) that run on GPUs have effectively been limited to CUDA—an extension of C/C++—or HIP from the AMD camp. Neither guarantees memory safety. Whether a pointer points to valid memory, and whether data races occur between threads, has always been entirely the programmer's responsibility.
This situation has weighed even more heavily since the 2020s, when GPUs became the core infrastructure for AI training and scientific computing. Bugs in kernels cause segmentation faults and undefined behavior, wasting days of computation on large clusters. NVIDIA's software ecosystem, with more than 250 CUDA libraries, has come to be called the "CUDA moat." Even when competing hardware appears, if software maturity fails to keep pace, migration simply does not happen.
Why Rust's ownership model never reached the GPU
Since its 1.0 release in 2015, Rust has grown its presence in systems programming as a language that guarantees memory safety and the elimination of data races through compile-time ownership checking. Adoption is progressing in the HPC (high-performance computing) community as well, including at Lawrence Livermore National Laboratory and the University of Toronto.
However, things were different when it came to GPU kernels specifically. Rust's ownership system is designed around the assumption that a single CPU thread accesses memory in a sequential order, and this is structurally at odds with the GPU's parallel execution model, in which thousands of threads run simultaneously. Existing attempts fell into three categories.
| Approach | Representative example | Limitation |
|---|---|---|
| SPIR-V-based compilation | rust-gpu | No support for generic pointers. Still transitioning from graphics to compute; lacks maturity |
| Thin bindings over vendor APIs | rust-cuda | Requires unsafe for every kernel. No memory-safety guarantee |
| Tight coupling to a specific vendor | cuda-oxide | NVIDIA-only. No portability |
A Rust GPU interface that was "safe, portable, and fast enough" did not exist.
Building on the common foundation of LLVM Offload
Filling this gap is the paper "GPU Offload in Rust: Portable, Safe, and Fast" (arXiv: 2608.13759) by a team of five led by Manuel S. Drehwald. Drehwald is pursuing a doctorate at the University of Toronto while working full-time on GPU offload development at Lawrence Livermore National Laboratory (LLNL). Co-author Johannes Doerfert is a principal developer of the LLVM Offload project, and Alán Aspuru-Guzik leads computational chemistry and AI research at the University of Toronto.
The core of their approach is to build GPU offload functionality directly into the Rust compiler (rustc) itself, generating native code for both NVIDIA and AMD GPUs via LLVM's Offload infrastructure. LLVM Offload was originally developed to support GPU execution for OpenMP, but was later decoupled from OpenMP and redesigned as a general-purpose foundation usable by any language frontend. By connecting Rust to this infrastructure—already proven with C++/Fortran—they opened a path that does not depend on vendor-specific toolchains.
Ownership information becomes "material" for compiler optimization
The technical heart of this work lies in the richness of the information that Rust's type system and ownership model supply to LLVM's intermediate representation (IR).
By default, safe Rust references are lowered into LLVM IR with $noalias$ metadata attached. $noalias$ is a declaration to the compiler stating: "no other pointer accesses the memory region that this pointer points to." In C/C++, this information can only be obtained if the programmer manually adds the restrict keyword; in Rust, it is attached automatically simply by writing safe code. The LLVM backend can use this information to perform optimizations more aggressively, such as reordering memory accesses and eliminating unnecessary loads.
Automatic generation of data transfers is also underpinned by the ownership model. At the MIR (Mid-level Intermediate Representation) stage, the compiler scans the types, layout, and mutability of kernel arguments to automatically determine what data, and how many bytes, should be transferred from host to device. Immutable references (&T) are sent to the device as read-only, while mutable references (&mut T) are treated as targets to be written back. The programmer does not need to explicitly write OpenMP's #pragma omp target map or CUDA's cudaMemcpy.
For safe parallel access inside GPU kernels, they introduced an abstraction called "Region." This is a mechanism for expressing, without breaking ownership rules, the typical pattern in which thousands of threads access different elements of the same slice. Internally, the frontend uses raw pointers, but the code the user writes remains safe Rust.
A head-to-head comparison against CUDA/HIP using RAJAPerf
For evaluation, they used the RAJAPerf benchmark suite developed by LLNL. RAJAPerf is a collection of loop-based computational kernels extracted from HPC applications, and performance comparisons against OpenMP, CUDA, and HIP implementations are standardized within it. Drehwald and colleagues ported a portion of this suite to pure Rust and took measurements across three environments: AMD MI250X, NVIDIA H100, and NVIDIA RTX A2000. The compiler used was an extended rustc based on LLVM 23.1.0-rc1.
Kernel execution time: roughly on par, with some advantages
In terms of standalone kernel execution time, Rust Offload showed performance roughly on par with RAJA (the C++ implementation). In two benchmarks, FIR and LTIMES, Rust was 44% and 46% slower than CUDA respectively, but these are tiny loops containing only a few multiply/add operations, where differences in the compiler's unrolling decisions directly affect the results. In the same two benchmarks, Rust was 15% and 32% faster than HIP (AMD). In other words, the team's analysis is that the slower cases leave room for improvement through fine-tuning of code generation.
Overall execution time: synchronization overhead is a challenge
When looking at overall execution time—from kernel launch through synchronization—the gap widens. On the MI250X, results ranged from 32% faster to 43% slower than CUDA/RAJA; on the H100, from 11% faster to 46% slower. The team acknowledges there is room for improvement in host-device synchronization efficiency.
Data transfer volume and transfer time
Summed across all benchmarks on the H100, Rust performed 53 host-to-device transfers totaling 423 MB, fewer than RAJA's 55 transfers totaling 468 MB. Nevertheless, the time spent on transfers was reversed: 46 ms for Rust versus 16 ms for RAJA. The team attributes this to differences in memory kind and to the fact that asynchronous transfers have not yet been implemented.
| Metric (H100, RAJAPerf overall) | Rust Offload | RAJA (CUDA C++) |
|---|---|---|
| Host→device transfer count | 53 | 55 |
| Host→device transfer volume | 423 MB | 468 MB |
| Host→device transfer time | 46 ms | 16 ms |
| Device→host transfer volume | 69 MB | 99 MB |
| Kernel execution time (near median) | Roughly on par | Baseline |
There is one more important figure. A naive implementation that repeats data transfer on every kernel launch (a simple application of Interface A) is more than 400 times slower on the MI250X than the optimized implementation (the absolute execution times for the compared and reference cases are shown only in the paper's figures, not stated in the text). To close this gap, the team extended LLVM's OpenMP-opt pass and built a prototype optimization that collectively eliminates repeated automatic transfers.
The structural reason safe code can become fast code
What these results signify is a reversal of the conventional wisdom that "safe means slow." Rust's ownership system is simultaneously a constraint on the programmer and a source of information for the compiler. Because $noalias$ guarantees are attached automatically, optimization opportunities that in C/C++ could only be obtained through manual annotation or profiling-based tuning become available simply by writing safe code.
Of course, this is not a claim that "Rust is always faster than CUDA." What the paper demonstrates is the fact that the quality of the IR generated by the LLVM backend has reached a level competitive with CUDA/HIP compilers, and there remain benchmarks where a gap persists. The team itself states that the slowdowns in FIR and LTIMES stem from differences in the compiler's unrolling decisions, and indicates that these gaps are expected to narrow through improvements in code generation.
Integration into rustc is underway, but the path to stable is still halfway
This research does not end with an academic publication alone. Drehwald is working in parallel on integrating GPU offload functionality into rustc itself, and a tracking issue (rust-lang/rust#131513) was opened in October 2024. In mid-2025, a PR for host-side code generation was merged, and progress has continued incrementally with kernel launches, simplification of the compilation process (ultimately reduced to two cargo commands plus one clang-linker-wrapper invocation), and enabling of testing in CI. "Completion of the std::offload module" was also adopted as a Rust Project Goal for the second half of 2025.
However, at present this remains an experimental feature on the nightly channel, and no timeline for stabilization has been announced. Intel GPU support is planned to be added once the LLVM-side backend matures, and there is also mention of Apple Silicon.
Remaining challenges and what this research is asking
There is no shortage of unresolved issues. First is the synchronization overhead between host and device. Even though kernel execution itself is on par, this is the primary factor behind the gap in overall execution time. Integrating asynchronous transfers with kernel launches is the next goal. Second, there is not yet a mechanism for running Rust's standard library (std) on the GPU. The team has indicated a plan to port LLVM's libc-for-gpu project for Rust, but implementation has not yet begun. Third, support for multi-device environments spanning multiple GPUs also remains at the planning stage.
On the matter of reproducibility, this paper is a preprint on arXiv and has not undergone peer review. The evaluation uses a subset of RAJAPerf and does not cover every kernel category.
Even so, the question this research raises touches on the fundamentals of GPU computing. Ever since 2006, writing code that runs on a GPU has meant giving up the guarantee of memory safety. The 20-year ecosystem CUDA has built rests on that trade-off. If kernels written in a safe language can hold their own in performance, the very design philosophy of GPU software could change. The next focal point is whether these results can be reproduced across a wider variety of applications and hardware configurations, and how far the synchronization overhead can be cut down before this lands in a stable release of rustc.