Fearless Concurrency on the GPU
Melih Elibol, Jared Roesch, Isaac Gelado, Eric Buehler, Michael Garland
Abstract
Rust has made safe systems programming practical on the CPU, but writing custom GPU kernels in Rust still forces programmers outside the language's ownership guarantees. We present cuTile Rust, a tile-based system for safe, idiomatic GPU kernel authoring in Rust. cuTile Rust extends Rust's ownership discipline to tile-based GPU kernels: mutable outputs are split into disjoint pieces, kernel launches preserve the host-side ownership contract, and programmers can opt out locally when they need lower-level control. The system also provides a composable host execution model spanning synchronous launches, asynchronous pipelines, and CUDA graph replay. Our evaluation shows that these abstractions can preserve performance on high-end GPUs. On the NVIDIA B200 GPU, cuTile Rust achieves 7 TB/s for element-wise operations and 2 PFlop/s for GEMM (96% of cuBLAS), matching cuTile Python within measurement noise. Grout, a cuTile-Rust-based inference engine, exercises cuTile Rust across an end-to-end Qwen3 inference path. In batch-1 decode, Grout reaches 171 generated tokens/s for Qwen3-4B on the NVIDIA GeForce RTX 5090 and 82 generated tokens/s for Qwen3-32B on the B200, competitive with vLLM and SGLang and consistent with an HBM roofline sanity check.
AI Impact Assessments
(1 models)Scientific Impact Assessment: "Fearless Concurrency on the GPU"
1. Core Contribution
This paper presents cuTile Rust, a system that extends Rust's ownership and borrowing guarantees to GPU kernel programming through a tile-based abstraction. The key insight is that the tile programming model creates a natural correspondence with Rust's aliasing XOR mutability rule: mutable output tensors are partitioned into disjoint sub-tensors (one per tile program), immutable inputs are broadcast as shared references, and Tile IR's token mechanism enforces intra-program ordering. The system spans the full host-device boundary with three interconnected contributions: (1) a safe kernel authoring model using partitioned mutable views and shared immutable views, (2) a generated launch interface that preserves ownership semantics across the CPU/GPU boundary, and (3) a composable `DeviceOp` execution model supporting sync, async, and CUDA graph replay modes.
The central problem being solved is that existing Rust GPU programming efforts (Rust-CUDA, rust-gpu, cuda-oxide) treat device code as `unsafe`, forfeiting the language's core safety guarantees. cuTile Rust recovers these guarantees for the class of kernels expressible in the tile model.
2. Methodological Rigor
The paper is technically well-constructed. The type grammar at the launch boundary (Figure 1) is precisely specified, and the host-to-device mapping (Table 1) is exhaustive. The data-race freedom proof (Appendix A) is clean and follows directly from two properties: alias-XOR-mutability across tile programs and token ordering within a tile program. The proof correctly reduces to the same argument structure that makes sequential Rust data-race-free.
The evaluation is structured around the right questions. The safety-overhead microbenchmarks compare safe Rust against unsafe Rust on the same backend (isolating the cost of safety abstractions), cuTile Python (isolating language-frontend effects), and cuBLAS (anchoring in absolute terms). The GEMM result—safe Rust matching unsafe Rust within 0.3% and reaching 96.4% of cuBLAS—is convincing evidence that the safety abstractions are zero-cost at steady state. The element-wise benchmark similarly shows no measurable overhead.
However, there are methodological gaps. The GEMM comparison only covers square matrices in powers of two; irregular shapes common in real workloads (e.g., non-square attention projections) are not evaluated. The end-to-end Grout evaluation is limited to batch-1 decode of a single model family (Qwen3), which is the easiest serving scenario. The paper acknowledges this but the inference evaluation would be stronger with batched scenarios or multiple model architectures.
3. Potential Impact
GPU programming safety: This work addresses a genuine pain point. As GPU-accelerated Rust applications grow (Candle, Burn, mistral.rs), the unsafe gap at the kernel boundary becomes increasingly consequential. Demonstrating that safety can be achieved at zero runtime cost removes the perceived performance tax that has historically deterred adoption of safer abstractions in GPU programming.
Tile-based programming: The paper contributes to the broader tile programming ecosystem (Triton, Pallas, ThunderKittens) by showing that tile abstractions naturally compose with ownership-based type systems. This could influence future language design for heterogeneous computing.
Practical adoption: The system is built on NVIDIA's Tile IR backend and developed at NVIDIA, suggesting potential for production deployment. The collaboration with Hugging Face on Grout signals industry interest. However, the NVIDIA-specific nature (CUDA, Tile IR) limits cross-platform applicability.
Composable execution model: The `DeviceOp` trait design—separating work construction from execution and supporting sync/async/graph through one interface—is a clean abstraction that could influence GPU runtime design beyond Rust.
4. Timeliness & Relevance
The paper is highly timely. Rust adoption in systems programming is accelerating, and the AI/ML ecosystem increasingly demands custom GPU kernels. The convergence of these trends creates demand for exactly this kind of system. The evaluation on modern hardware (B200, RTX 5090) and current models (Qwen3) grounds the work in today's workloads.
The tile programming model is also timely, with Triton having demonstrated the productivity benefits and NVIDIA's own investment in Tile IR signaling architectural commitment. Bringing Rust's safety guarantees to this model is a natural and well-timed contribution.
5. Strengths & Limitations
Key Strengths:
Notable Limitations:
Additional Observations
The branded `PartitionIndex` mechanism for mapped partitions is an elegant use of Rust's type system to carry disjointness proofs without runtime checks. The paper could have emphasized this contribution more, as it solves a non-trivial problem (multi-output-per-program safety) that goes beyond simple partitioning.
The paper is well-written and clearly structured, though dense. The code examples effectively illustrate the programming model's feel.
Generated Jun 16, 2026
Comparison History (39)
Paper 1 offers a highly rigorous evaluation on state-of-the-art hardware, demonstrating that Rust's safety guarantees can be applied to GPU programming without sacrificing performance (matching cuBLAS). This solves a critical, widespread problem in high-performance computing and AI inference. While Paper 2 addresses a timely topic (LLM formal math), its results rely on 'early experiments' and lack the methodological maturity and immediate, broad real-world applicability demonstrated by Paper 1's end-to-end inference engine.
Paper 2 presents a breakthrough in applying Rust's safety guarantees to high-performance GPU programming, a critical challenge in modern systems engineering. Its ability to achieve near-native cuBLAS performance and competitive LLM inference speeds demonstrates immense real-world applicability. While Paper 1 offers valuable compiler infrastructure for the niche field of neuromorphic computing, Paper 2 tackles a pervasive safety and performance problem in mainstream AI and accelerated computing, giving it significantly higher breadth of impact, timeliness, and potential for widespread adoption.
Paper 1 introduces a novel system extending Rust's ownership model to GPU programming, addressing a fundamental gap in safe GPU kernel authoring. It demonstrates near-optimal performance (96% of cuBLAS) on cutting-edge hardware, with practical applications in LLM inference. The combination of memory safety guarantees with high-performance GPU computing has broad implications across HPC, ML, and systems programming. Paper 2 makes a solid but incremental contribution to formal verification of embedded C code with limited scope (two case studies in automotive). Paper 1's timeliness (GPU computing + Rust + LLMs) and breadth of impact are significantly greater.
Paper 2 addresses the highly timely intersection of Rust's safety guarantees and GPU programming, with immediate practical applications in AI/ML inference. It demonstrates near-optimal performance (96% of cuBLAS) while maintaining memory safety, which could broadly impact GPU programming practices. Paper 1 contributes meaningfully to probabilistic program verification but targets a more niche community. Paper 2's breadth of impact across systems programming, GPU computing, and AI deployment, combined with its practical demonstration on current hardware and competitive LLM inference results, gives it higher potential scientific impact.
Paper 1 likely has higher impact: it introduces a novel, rigorous programming-model extension (ownership/safety guarantees) for GPU kernels in Rust, addressing a long-standing correctness/productivity gap with broad relevance to PL, systems, and HPC/AI infrastructure. It demonstrates strong methodological rigor with performance near hardware/roofline and cuBLAS, plus an end-to-end inference engine validation. Real-world applicability is immediate for safe high-performance GPU software and could influence language and GPU programming ecosystems. Paper 2 is timely and useful, but its MoE specialization is more incremental within a crowded LLM codegen space and narrower in cross-field impact.
Paper 1 addresses a high-impact problem at the intersection of systems programming, GPU computing, and AI inference—areas of enormous current interest. Bringing Rust's safety guarantees to GPU programming is novel and practically significant, with demonstrated near-peak performance (96% of cuBLAS) and a working LLM inference engine. Its breadth of impact spans systems programming, HPC, and ML infrastructure. Paper 2 makes solid theoretical contributions to program synthesis for nonlinear real arithmetic, but addresses a more niche problem with narrower applicability. Paper 1's timeliness given the GPU computing boom gives it the edge.
Paper 1 presents a novel system extending Rust's ownership model to GPU programming, addressing a significant gap in safe systems programming for accelerators. It demonstrates near-peak performance (96% of cuBLAS) while maintaining safety guarantees, and validates with an end-to-end LLM inference engine. This combines high novelty (first practical safe GPU kernel authoring in Rust), immediate real-world applicability (GPU/AI workloads), and timeliness (AI infrastructure is critical). Paper 2 addresses compiler phase ordering with ML, which is a well-studied problem with incremental advances, offering less transformative impact despite solid results.
Paper 1 presents a highly practical and timely solution for safe GPU programming using Rust, demonstrating competitive performance for modern LLM inference on high-end GPUs. Its immediate applicability to AI systems and GPU computing gives it significantly broader potential real-world impact and relevance compared to Paper 2, which focuses on niche theoretical complexity bounds for a fragment of Datalog.
Paper 2 likely has higher impact: it introduces a safe, idiomatic Rust model for GPU kernel programming with strong real-world applicability to ML/HPC, demonstrating near–state-of-the-art performance (e.g., 96% cuBLAS GEMM) and an end-to-end LLM inference engine. This combines novelty (extending ownership/borrow checking to GPU tiles and launches), methodological rigor (performance/roofline validation), timeliness (GPU programming + Rust + LLM inference), and broad cross-field relevance (systems, compilers, PL, ML, HPC). Paper 1 is strong but more specialized to graph pattern mining optimizers.
Paper 2 likely has higher impact due to strong timeliness (GPU programming and LLM inference), clear real-world applicability, and compelling evidence of practicality and performance (near-cuBLAS GEMM, end-to-end inference results competitive with major systems). Extending Rust’s safety/ownership model to GPU kernels is a novel, broadly relevant contribution spanning PL, systems, and ML infrastructure. Paper 1 is innovative within compiler optimization research, but appears more prototype-stage and narrower in immediate deployment scope, with impact dependent on broader adoption and demonstrated wins on real compiler workloads.