In 2006, when NVIDIA released CUDA (Compute Unified Device Architecture), GPUs were transformed from components dedicated solely to graphics rendering into general-purpose parallel computers. Since then, the practical means of writing kernels (functions that carry out the core of a computation) that run on GPUs have effectively been limited to CUDA, an extension of C/C++, or HIP on the AMD side. Neither guarantees memory safety. Whether the memory a pointer points to is valid, and whether data races occur between threads, are entirely the programmer's responsibility.
This situation has only weighed more heavily as GPUs became the central infrastructure for AI training and scientific computing in the 2020s. Bugs in kernels cause segmentation faults and undefined behavior, wasting days of computation on large clusters. NVIDIA's software ecosystem holds over 250 CUDA libraries and has been called "the CUDA moat." Even when competing hardware appears, switching does not happen unless software maturity keeps pace.
Why Rust's ownership model never reached the GPU
Since its 1.0 release in 2015, Rust has grown in prominence in systems programming as a language that guarantees memory safety and eliminates data races through compile-time ownership checks. Adoption is also progressing in the HPC (high-performance computing) community, including at Lawrence Livermore National Laboratory and the University of Toronto.
However, when it comes to GPU kernels specifically, the situation was different. Rust's ownership system was designed on the premise that a single CPU thread accesses memory in a sequential order, and this is structurally incompatible with the GPU's parallel execution model, in which thousands of threads run simultaneously. Existing attempts fell into three directions.
| Approach | Representative example | Constraint |
|---|---|---|
| SPIR-V based compilation | rust-gpu | No support for general-purpose pointers. Still transitioning from graphics to compute, and lacks maturity |
| Thin bindings to vendor APIs | rust-cuda | unsafe required for every kernel. No memory safety guarantee |
| Tight coupling to a specific vendor | cuda-oxide | NVIDIA-only. No portability |
A "safe, portable, and fast enough" Rust GPU interface simply did not exist.
Building on the common LLVM Offload infrastructure
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 a PhD student 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 lead 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 infrastructure that any language frontend could use. By connecting Rust to this infrastructure, which already has a track record with C++/Fortran, the team opened a path that does not depend on vendor-specific toolchains.
Ownership information as "material" for compiler optimization
The technical crux lies in the richness of the information that Rust's type system and ownership model provide to LLVM's intermediate representation (IR).
References in safe Rust are lowered into LLVM IR with $noalias$ metadata by default. $noalias$ is a declaration to the compiler that "no other pointer will access the memory region this pointer points to." In C/C++, this information is only available if the programmer manually attaches the restrict keyword; in Rust, it is attached automatically simply by writing safe code. The LLVM backend can use this information to more aggressively perform optimizations such as reordering memory accesses and eliminating unnecessary loads.
Automatic generation of data transfers is also supported by the ownership model. The compiler scans the types, layout, and mutability of kernel arguments at Rust's MIR (Mid-level Intermediate Representation) stage, and automatically determines what data, and how many bytes, should be transferred from the host to the 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. Programmers do not need to explicitly write OpenMP's #pragma omp target map or CUDA's cudaMemcpy.
For safe parallel access inside GPU kernels, the team introduced an abstraction called "Region." This is a mechanism for expressing the typical pattern in which thousands of threads access different elements of the same slice, without breaking the rules of ownership. Internally, the frontend uses raw pointers, but the code the user writes remains safe Rust.
Head-to-head comparison with CUDA/HIP using RAJAPerf
For evaluation, the team 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 part of this suite to pure Rust and measured performance 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: nearly equivalent, superior in some cases
In terms of standalone kernel execution time, Rust Offload showed performance nearly equivalent to RAJA (the C++ implementation). On two benchmarks, FIR and LTIMES, Rust was 44% and 46% slower than CUDA respectively, but these are tiny loops containing only a few multiplications/additions, so differences in the compiler's unrolling decisions directly affect the results. On the same two benchmarks, Rust was 15% and 32% faster than HIP (AMD), respectively. In other words, the team's analysis is that the slower cases have room to be closed through fine-tuning of code generation.
Overall execution time: synchronization overhead remains a challenge
The gap widens when looking at overall execution time, from kernel launch to synchronization. On the MI250X, results ranged from 32% faster to 43% slower compared to CUDA/RAJA; on the H100, from 11% faster to 46% slower. The team acknowledges that there is room for improvement in host-device synchronization efficiency.
Data transfer volume and transfer time
Across all benchmarks combined on the H100, Rust performed 53 host-to-device transfers totaling 423 MB. This is less than RAJA's 55 transfers totaling 468 MB. Nevertheless, the transfer time 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 transfer has 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) | Nearly equivalent | Baseline |
There is another 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 compared to the optimized implementation (the absolute execution times being compared are shown only in the paper's figures and are not stated in the text). To close this gap, the team extended LLVM's OpenMP-opt pass and built a prototype optimization that batches and eliminates repeated automatic transfers.
The structural reason why 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, at the same time as being a constraint on the programmer, also 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 are now 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 prospects for narrowing this through improvements in code generation.
Integration into rustc is underway, but a stable release is still a ways off
This research does not end with an academic publication alone. Drehwald is concurrently working to integrate 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 step by step on kernel launching, simplification of the compilation procedure (ultimately reduced to two cargo commands and one clang-linker-wrapper invocation), and enabling tests in CI. The completion of the std::offload module was also adopted as one of the Rust Project Goals for the latter half of 2025.
However, at present this remains an experimental feature on the nightly channel, and no timeline for landing in a stable release has been specified. Intel GPU support is planned to be added once the backend on the LLVM side matures, and there is also mention of Apple Silicon.
Remaining challenges, and the question this research raises
There is no shortage of unresolved issues. First is the synchronization overhead between host and device. Even though kernel execution itself is comparable, this is the main cause of the gap seen in overall execution time. Integrating asynchronous transfer with kernel launch 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 policy of porting LLVM's libc-for-gpu project for Rust, but the implementation is still to come. Third, support for multi-device environments spanning multiple GPUs also remains at the planning stage.
In terms of reproducibility, this paper is a preprint on arXiv and has not undergone peer review. The evaluation used a subset of RAJAPerf and did not cover all kernel categories.
Even so, the question this research raises touches on the very foundations of GPU computing. Since 2006, writing code that runs on a GPU has meant giving up memory safety guarantees. The 20-year ecosystem that CUDA has built rests on top of that trade-off. If kernels written in a safe language turn out to be comparable in performance, the very design philosophy of GPU software could change. The next focal points are whether these results can be reproduced across a wider variety of applications and hardware configurations, and how much of the synchronization overhead can be trimmed away before this lands in a stable release of rustc.