编译器放行了，数据竞争还在：CUDA Rust 的两条路，交易的是控制权

想象一个每天都在发生的场景。你把手写 CUDA kernel 用 Rust 重写了一遍，rustc 干净利落地放行，没有告警，没有 unsafe 报错。你会下意识地松一口气：类型系统在看着，数据竞争应该没有了。但如果你真的把这段代码丢到 GPU 上跑，两个 block 可能同时往同一段显存里写，结果还是错的。这不是 Rust 的 bug，而是它能力边界的一部分：借用检查器能保证的，是语言内部的引用别名安全，&mut 是否独占、生命周期是否越界，这些结论只靠类型本身就能证明。而 GPU 上的数据竞争发生在类型系统看不见的地方：一次 launch 把成百上千个线程送进同一个函数体，一份 &mut 切片被所有线程同时握着，线程的身份、数量、访问哪一段内存，编译器在类型层无从得知。所以编译通过从来不等价于没有数据竞争，"换成 Rust 就自动消除 GPU 数据竞争"这个直觉，错就错在把语言的证明范围当成了 GPU 的证明范围。那么问题就移到了语言之外：这个证明由谁来补？NVIDIA 官方给出的两条 Rust 路径，cuda-oxide（SIMT 路径）和 cutile-rs（Tile 路径），正是对这个问题的两种后续解答。接下来我们就沿着这两条路往下走，看安全证明分别被写进了哪里，代价又落在谁身上。

1 「换成 Rust 就安全了」，这个直觉错在哪

先回到 Rust 的安全承诺本身。它防数据竞争的机制，核心是借用检查器（borrow checker）加上两个标记 trait：Send 和 Sync。Rust 官方文档 The Rustonomicon 给的定义很简短：一个类型是 Send，指的是把它的值送到另一个线程去用是安全的；一个类型是 Sync，指的是它可以在线程之间共享，等价写法是 &T 本身是 Send。这两个标记是自动推导的，一个复合类型只要所有成员都 Send，它自己也 Send。绝大多数标准库类型都自动满足，所以三个显式例外格外醒目：裸指针（raw pointer）既不是 Send 也不是 Sync，因为它没有任何安全护栏；UnsafeCell 不是 Sync，于是基于它的 Cell、RefCell 也不是；Rc 不是 Send 也不是 Sync，因为它的引用计数是共享且无同步的。Rustonomicon

看下面这张类型矩阵，一眼就能分辨哪些标准库类型自动跨线程（绿色），哪些必须显式证明（红色）：

图注：Send/Sync 类型矩阵。&mut T 一格是绿色的，因为它是 Send。但注意底部引用框里的前提：这个标记只说明"类型层面没有内部可变性、没有 Rc、没有裸指针可达"，它完全没有提到 GPU kernel 的网格几何。Rustonomicon

问题就出在这里。借用检查器只理解语言内的引用别名，它认识 &mut、认识生命周期，但它不认识线程。确切地说，它不认识 GPU 的启动几何（launch geometry）：网格多大、每个 block 多少线程、每个线程写哪一段内存，这些信息在 Rust 语言里根本没有表征。&mut [T] 在类型层确实是 Send，这个结论只靠类型本身就能证明，编译器不需要知道"两个线程拿到的分区互不重叠"。所以在 CPU 上，thread::spawn(move || ...) 把切片移进新线程，编译器可以确认原线程已经交出了访问权。但在 GPU 上，一次 launch 把成百上千个线程送进同一个函数体，每个线程手里都握着同一份 &mut 指针，编译器却把它当成"唯一的写者"看。它检查的是签名，不知道持有签名的人是谁。

借用检查器就像只检查门禁卡上有没有写"允许进入"，至于进来的是一个人还是一百个人，它看不见。

这里要补一句学术背景，免得"谁补这个证明"听起来像玄学。2018 年 POPL 上的 RustBelt 论文（MPI-SWS 与 TU Delft，与 NVIDIA 无关）用 λRust 和 Iris 并发分离逻辑，在 Coq 里给 Rust 类型系统做了机器可检验的语义证明，用 ownership predicate 和 sharing predicate 两条轨道解释内部可变性。它证明了一件事：unsafe 封装的安全边界是可以被严肃形式化讨论的，不是靠"我觉得应该没问题"。RustBelt

所以第一层的结论是：Rust 在语言内的别名证明是可靠的，但 GPU 的线程身份在类型系统之外。编译放行从来不等价于没有数据竞争。真正的问题是下一个：既然语言自己补不了这个证明，类型系统本身能不能补？

2 SIMT 路径：把「一线程一输出」写进类型，把控制权的代价留给 unsafe

能，而且 NVIDIA 的 cuda-oxide 就是这么做的。它走的是 SIMT 路径：保持你熟悉的线程模型，但把"每个线程独占一片输出"这件事从文档约定变成类型约束。

先看它的入口代码长什么样。cuda-oxide 的 vecadd 示例是这么写的：

图注：cuda-oxide 的 SIMT kernel 入口。三行属性宏声明了这是 kernel、block 上限 256 线程、以及 #[launch_contract(domain = 1, block = (256, 1, 1))] 声明的启动契约。注意 c 参数的类型不是 &mut [f32]，而是 DisjointSlice<f32>，索引来自 thread::index_1d() 而不是裸 usize。NVIDIA Blog

核心是 DisjointSlice<T, IndexSpace> 这个类型。直译是"不相交切片"，它把"互不重叠的内存分区"直接写进了类型签名。配合它的是 ThreadIndex<'kernel, IndexSpace>，一个带 'kernel 生命周期、且 !Send + !Sync + !Copy + !Clone 的不透明句柄，术语叫 witness（见证）。这两件套合起来的效果是：你想写 c.get_mut(idx)，必须拿一个 ThreadIndex 来交换，而这个 witness 只有 thread::index_1d() 这类受信任的构造器能造出来，线程之间无法借共享内存偷渡。IndexSpace 这个类型参数还负责把 witness 钉死在特定布局上，Index2D<256> 的 witness 喂给 Index2D<128> 的分区，编译期直接拒绝。

这就是把"一线程一输出"证明进了类型层：写入互不重叠这件事，不再依赖程序员自律，而是类型系统强制。cuda-oxide

但类型化的切片只是半边。另一半是启动契约（launch contract）。#[launch_contract(domain = 1, block = (256, 1, 1))] 声明了网格几何之后，host 端必须走 prepare_* 系列函数来准备 launch，而不是直接用裸的配置。这一套检查会在 host 端核对 block 形状与契约完全一致、设备限制、动态共享内存、cluster 与 cooperative 约束；设备端编译器则把契约译成 .reqntid 指令，让 CUDA driver 拒绝任何形状不一致的 launch。两道闸门把"几何写错"和"内核错配"都挡在门外。下面这张图展示的是完整的 launch 生命周期：

图注：host 端六步 launch 流程（Init context → Load module → Get function → Configure → Launch → Synchronize）与设备端 SM 调度。注意红色虚线框里的 cuda_launch! 是 unsafe 逃生口，安全的 typed launch 走的是 prepare_* + PreparedLaunch 那条链，不在这张图的流程里。cuda-oxide

这一整套设计里，编译链是值得一提的。cuda-oxide 不是把 Rust 翻译成 CUDA C 再编译，而是真的走 rustc frontend，从 Stable MIR 进入自己基于 Pliron 的 IR，最后落到 LLVM NVPTX 后端生成 PTX：

图注：cuda-oxide 编译链的七个节点：Rust source → rustc frontend → Stable MIR → MIR dialect（pliron）→ LLVM dialect（pliron）→ LLVM IR → PTX，右侧大括号归入 rustc、pliron、LLVM 三个阶段。安全证明发生在类型层与 host 验证，不是换了个编译器就自动有的。cuda-oxide

现在到了这条路径最关键的分寸感。cuda-oxide 的 Safety Model 把安全分成三层：Tier 1 是可以编译期证明的部分，DisjointSlice 加 PreparedLaunch 就在这一层；Tier 2 是 unsafe 但有契约（contract）的部分，共享内存、动态共享内存、shfl_sync 这类 warp 原语、barrier，都在这层，使用它们有明确的安全约定；Tier 3 是手工硬件 intrinsic，TMA、tcgen05、WGMMA、Cluster 这类直接摸硬件的操作。也就是说，cuda-oxide 自己文档里承认的 NOT enforced 清单有四项：&mut [T] 作 kernel 参数（官方处置是"凡收这种参数的内核，就当每一行都是 unsafe"）、thread divergence（rustc 的跳转线程化优化会破坏 bar.sync 的整块汇合，现在只能把这个 pass 关掉，是 disable 不是 fix）、warp 级收敛（shfl_sync 要求 mask 内线程在调用点汇合，类型系统管不了，出错的形态是静默挂起）、memory-space awareness（global、shared、TMEM 对借用检查器不可见，只能一律保守对待）。

"控制越细，证明债务越重"，这句话在 SIMT 路径上被 NOT enforced 四项钉死了。你想要的 warp 协作、共享内存吞吐、收敛语义，恰恰都是高性能 SIMT kernel 的命根子，而它们全留在 unsafe 里。Fearless Concurrency on the GPU 论文在 Related Work 里也把这一类 SIMT kernel 列为"仍被视为 unsafe"的典型。Fearless GPU

那自然的下一问是：如果放弃"每个线程都由程序员安排"这个前提，把线程映射整个交给编译器，证明债务能省到什么程度？

3 Tile 路径：把线程映射交给编译器，所有权跨过 launch 边界

这就是第二条路，cutile-rs 背后的思路。Tile 路径做了一个 SIMT 做不到的抽象跃迁：程序里根本没有"线程"这个概念。你声明的是对整块张量的运算，至于这块运算怎么分配到线程上，由编译器决定。

机制的核心是张量分区（partition）。host 端拿到一个 &mut Tensor<T, S> 表示可变输出，launch 之前必须调用 partition([B]) 把它切成不重叠的子张量。每个 tile program（可以理解成 GPU 端的一个逻辑执行单元）独占一个分区视图（partition view），而只读输入 &Tensor<T, S> 则广播共享给所有 tile program。看论文里的这张图：

图注：host 到 device 的分区映射。Partition<Tensor<T>> 把可变张量切成多个子张量，每个 tile program 拿到一个独占的 &mut 子张量（粉色，标注 exclusive 1:1）；青色长条带是 &Tensor<T> 的广播共享。注意这里的单元是 tile program，不是线程。Fearless GPU

对照第 2 章你就看到区别了。SIMT 路径要程序员手工写"这个线程写这个位置"，Tile 路径里连这个念头都没有：每个 tile program 拿到的分区天然互不重叠，互斥性由 partition API 和借用检查器在类型层保证，错误根本没有表达的空间。论文举了一个很直观的例子：cuTile Python 里如果把 store 的目标下标 m 和 h2 写反，在 grid (B·H, H, 1) 下会有 17% 到 35% 的元素输出不确定；同样的错误在 cuTile Rust 里根本无法写出，因为 store 的目标不是程序员选的 index，而是 partition view 本身。

不止分区间不重叠，同一分区内的操作顺序也有类型约束。cutile-rs 的 host 端 API 把 GPU 工作描述成 lazy 组合的 DeviceOp，&mut 视图上的每个 load/store 会被 token 链串起来，构成 happens-before；只读共享输入则没有 token 约束，编译器可以自由重排：

图注：单个 tile program 内的 token 排序。左侧粉色框 Tᵢ 是该 program 的可变子张量，t₀ → t₁ → t₂ 的 token 链把 c.load()、c.store() 串成 happens-before；右侧青色共享输入与 Tᵢ 之间是虚线的 no token constraints，编译器自由重排。token 是 Tile IR 的 happens-before 边，不是 warp barrier。Fearless GPU

这两条事实合起来，就是论文 Appendix A 那条形式化定理的形象版：safe 的 cuTile Rust kernel 在 Tile IR 内存模型下数据 race-free。证明的思路很干净，分两种情况。跨 tile program 的冲突对，因为写者只能通过自己的 partition view 写，而分区互不重叠，矛盾。同 program 内的冲突对，因为写只走这个 program 的 &mut 视图，所有访问在 program order 和 token order 双重相关。都构不成 data race。

跨过 launch 边界的所有权延续，是靠 prepare/recover 协议实现的，仓库文档里把它表述成"你交进去什么，就拿回来什么"（same-type-in-same-type-out）：Arc<Tensor> 进、Arc<Tensor> 出，分区进、分区出。只读输入 &Tensor<T> 不是 'static，tokio::spawn 在编译期就会拒绝借用形式的操作，连运行时检查都不需要。

这个抽象层不是 Rust 凭空发明的，它上面有一个正式的落点，叫 Tile IR。NVIDIA 在 2025 年 12 月随 CUDA 13.1 发布了 Tile IR，官方定位是"为 tile 程序设计的虚拟指令集"，并且直接说它之于 tile 编程，类似于 PTX 之于 SIMT。cutile-rs 只是它众多前端里的一个 Rust 前端，其他前端还有 cuTile Python、cuTile.jl、Triton 后端。CUDA Tile IR

软件栈里的位置看这张图：

图注：NVIDIA 软件栈中两条编译路径并列。左列灰色是 SIMT path（CUDA C++/Python → NVVM/LLVM → PTX），右列绿色是 Tile path（cuTile → Tile IR），两条路径最终汇入同一块 GPU 芯片。两条路径是共存关系，不是取代关系。CUDA Tile IR

Tile 路径里，线程、内存层级、Tensor Core 的映射全部由 Tile IR 编译器接管，源码不再编码架构特定选择，这就是它跨代适配（sm_80 到 sm_100+）的来源。张量数据流本身还是传统的那一套，HBM 里的 Tensor 加载成寄存器里的 Tile，算完再存回去：

图注：Tile 数据流。input tensor（HBM 全局内存）→ load_tile → tile（寄存器）→ compute → result → store → output tensor，中央绿色高亮"Many tiles process concurrently"，分区后每个 tile 独立走这条 Load → Compute → Store 流。cutile-rs

但控制权不会凭空消失，它只是换了存放的地方。交出线程映射的代价是：warp 协作、共享内存、TMA 这些小颗粒硬件的直接入口都没有了，想要就得走 unsafe。cutile-rs 明确留了两条 opt-out：unsafe push_device_ptr 用于和裸 CUDA 互操作，unsafe fn then_unchecked 绕过"同一线程同时只有一个 DeviceOp"的执行锁。框架的立场是，#[cutile::entry] 生成的 launcher 是安全的，因为两端都在框架管理之下；你要自己伸手进去，就自己负责。cutile-rs

安全收益来自抽象层抬升：编译器接管线程映射，可证明面积随之扩大。代价是你失去了 warp 协作与共享内存的直接入口。那么剩下的悬念就一个了：交出这么多控制权，性能税有多大？

4 抽象收费吗：B200 实测几乎免费，证据却有边界条件

直觉上，包了一层类型系统又包了一层抽象层的代码，应该比手写的慢。Fearless Concurrency on the GPU 论文在 B200 上的实测给出的答案是：规则张量负载上，几乎不收费，且这个结论可以看得很细。

先看元素操作。N = 2²⁸ 的 element-wise add 是内存受限负载，safe Rust 达到 7.02 TB/s，cuTile Python 达到 7.01 TB/s，而 B200 的理论 DRAM 峰值是 7.68 TB/s，三者在一条线上。再看 GEMM，M=N=K=8192、f16，safe Rust 达到 2.07 PFlop/s，是 dense f16 峰值的 92%，是 cuBLAS 的 96.4%。注意这里有一个容易写反的数字：unsafe Rust 在同一份 schedule 下与 safe Rust 的差距在 0.3% 以内，这是"安全机制本身是否有成本"的内部对照，不是与底层 Tile IR 的差距。安全机制在 Tile IR 这一层确实接近零成本。Fearless GPU

三联子图把这条证据摆得很直观：

图注：安全开销微基准的三联子图。(a) element-wise 带宽，三组点贴近 7.68 TB/s 峰值虚线；(b) GEMM，safe Rust 顶部接近 2 PFlop/s，cuBLAS 同高，unsafe 与 safe 几乎重合，cuTile Python 略低；(c) JIT 开销是首次调用的编译成本，不是稳态成本。这是微基准，不是推理端到端结果。Fearless GPU

论文还报告了端到端的推理数字：Grout 引擎（用 cuTile Rust 写的 Qwen3 推理引擎）在 RTX 5090 上跑 Qwen3-4B 达到 171 tok/s，在 B200 上跑 Qwen3-32B 达到 82 tok/s。但这两个数字有一个必须说清的边界：它们是 batch-1 decode、单请求的实验结果，论文自己声明并发 batching 下的吞吐是另一项研究。把 171 当成"多用户并发下也这样"，就超出了证据范围。

而且，证据在反方向也有。Grout 仓库自己提交的 B200 benchmark bundle 显示：短生成的序列下（tg 从 36 到 8192），grout 以 79.6 到 80.1 tok/s 领先 vLLM（78.8 到 79.2）和 SGLang（74.4 到 77.4）；但把 prompt 拉到 8192 token 的长语境（pp=8192）时，grout 掉到 37.2 tok/s，反而被 vLLM 的 37.3 压过。同一个引擎，同一个 backend，换个负载形状排名就翻转，这是"性能数字必须绑定负载类型"最诚实的注脚。Grout

"也不是一定免费"的证据来自不规则负载。有一篇独立论文把同一个 hash-blocked 三维重建内核用 CUDA C++、Rust（cuda-oxide）和 Triton 三套实现做了实测。规则更新阶段三栈差距不大，但到了不规则分配阶段，Triton 与 CUDA C++ 的差距拉到 11.2 到 31.6 倍。更严重的是正确性问题，Triton 在 hash load factor 达到 0.071 起就开始静默丢块。这不是跑得慢，而是产出错误结果。Irregularity

空洞的形状能说明问题的性质：

图注：三维重建表面的"沉默"。图中央整块缺失的空洞，是"CUDA C++ 有而 Triton 没有"的体素点。损失不是散落的噪声，而是一整块 8³ 体素连着表面一起消失，因为 Triton 的 probe 循环必须绑死在编译期上限，探不到就整块丢。Irregularity

这篇论文跑的平台是 RTX 5070 Ti 和 RTX 5060，不是 B200，所以具体比值不能直接外推；但它的方法论结论，也就是高层抽象在不规则负载上可能同时付出性能和正确性双重税，与硬件无关。这类负载恰好是三维重建、粒子模拟这些真实生产场景的常态，而不是 GEMM 那样被所有工具链优化到极致的花园。

最后还要提一条和 NVIDIA 无关的路线，它决定"抽象税"这个问题不能被读窄。GPU Offload in Rust 这篇论文（作者来自 LLNL 与多伦多大学）走的是 rustc 加 LLVM Offload 的三遍流水线，目标是跨厂商：一份 Rust 源码可以同时打包 NVIDIA 和 AMD 的设备镜像，在 RAJAPerf 子集上与手调 CUDA/HIP 基线相比落在 ±45% 区间内。也就是说，"在 NVIDIA 生态内换抽象层"和"跨厂商移植"是两条不同的路，前者有 B200 实测背书，后者也有自己的实测，只是数字范围宽得多。GPU Offload

图注：rustc + LLVM Offload 的三遍实现流水线：Pass 1 在 host 收集 monomorphization 清单，Pass 2 在 device 侧编译并打包设备镜像，Pass 3 在 host 嵌入并链接出 a.out。图中带 [Automated] 与 [Automation in progress] 阴影的节点是进行中状态，不是生产可用。GPU Offload

所以第 4 章的结论是双面的：抽象税在规则张量上接近零，这是 B200 实测支持的；但不规则负载会让证据失效，长上下文会让排位翻转。任何"抽象不收费"的主张都必须绑定负载类型和 batch 条件。那问题就到了最后一步：证据都有边界，到底怎么选？

5 怎么选：按证明责任挑抽象层，可移植性止于 NVIDIA

NVIDIA 自己的博客给了一个非常明确的推荐顺序：优先 Tile。原文说得很直接："当你选一个来构建时，先伸手够 Tile。编译器决定 tile 如何映射到每个架构，所以你的源码不编码架构特定选择；需要那份控制权、想自己管理内存和线程时，再降到 SIMT。" 同时官方也给出了成熟度边界："两个项目都处于早期阶段，都还没有生产就绪。cuda-oxide 是 early alpha，cutile-rs 走得稍远，已经发布到 crates.io，并被 HuggingFace 的 Grout 和 mistral.rs 采用。"NVIDIA Blog

工具链门槛也支持这个顺序。cuda-oxide 要 pinned nightly（仓库当前钉在 nightly-2026-08-28）、CUDA 13.0+、LLVM 21+，NVIDIA 自己都承认这是想消除的摩擦；cutile-rs 只要 stable Rust 1.89+ 和 CUDA 13.3，不用 nightly、不用自带 LLVM。两个门槛放在一起，选 Tile 的门槛明显更低。

但"优先 Tile"不是无条件的。把第 2、3 章的 unsafe 材料放在一起看，选型其实是一张按证明责任对号的表。规则张量计算（GEMM、element-wise、attention 这类形状规整的 load/store）交给 Tile，编译器接管线程映射，可证明面积最大，性能税接近零，Grout 就是证例。依赖共享内存、warp 原语或者需要极致手工调优的内核（flash attention 的手写变体、各类融合 kernel），留在 SIMT，自己背 NOT enforced 四项的证明债务。想拿回控制权的时刻，两条路都留了 unsafe 出口：Tile 侧包括 raw pointer kernel entry（DevicePointer<T>）、unsafe push_device_ptr 与 unsafe fn then_unchecked，SIMT 是 Tier 2 契约和 Tier 3 硬件 intrinsic。Rust 的现实价值从来不是消灭 unsafe，而是把 unsafe 压缩成可审计的边界。这个思路不是新发明，RustBelt 在 2018 年就已经为之提供了形式语义的根基。RustBelt

最后是"可移植"这个词要拆开。NVIDIA 说的可移植，是架构可移植：同一份源码在 sm_80 到 sm_100+ 各代 NVIDIA 架构间不用重写。cutile-rs 的支持矩阵明确是 sm_80 起步，sm_90 需要 CUDA 13.3，sm_70、sm_75 不在路线图上。cutile-rs NVIDIA 官方把语言互操作（CUDA Rust 与 CUDA C++、CUDA Python 互通）当作路线图，但这是 NVIDIA 生态内的互操作，不是跨厂商抽象。真正跨厂商的是第 4 章那张三遍流水线代表的路线，rustc 加 LLVM Offload，机构是 LLNL 与多伦多大学，与 NVIDIA 无关。架构可移植不等于跨厂商可移植，把这两件事混为一谈，是 NVIDIA 官方叙事里最容易让人误读的地方。

把整条线收回来。开头那个幻觉是："换成 Rust 就安全了"。现在可以给出完整的修正版判断：Rust 没有自动让 CUDA 安全，它只是把"证明"这件事变成了一道选择题。想要对线程和内存的精细控制，就留在 SIMT，同时接受 NOT enforced 四项那样沉甸甸的 unsafe 债务；想要更大的可证明面积，就把线程映射交给编译器，走 Tile，同时接受失去 warp 协作和共享内存直接入口，以及任何"可移植"都止于 NVIDIA 生态这个事实。两条路交易的本质是控制权，谁掌握线程与内存映射，谁就承担正确性证明。这个判断框架，比任何"Rust 胜过 CUDA C++"的结论都更能带到明天的工作里。

参考文献

本专题引用的核心资料如下，建议深入学习时阅读原文：

1. Fearless Concurrency on the GPU — arXiv:2606.15991
2. Introducing CUDA Rust: Two Tracks for Writing GPU Kernels — NVIDIA Developer Blog
3. NVlabs/cuda-oxide & The cuda-oxide Book — GitHub / Book
4. NVlabs/cutile-rs & The cuTile Rust Book — GitHub / Book
5. huggingface/grout: Qwen3 inference engine on cuTile Rust — GitHub
6. GPU Offload in Rust: Portable, Safe, and Fast — arXiv:2608.13759
7. What Irregularity Costs: CUDA C++, Rust, and Triton on a Hash-Blocked GPU Workload — arXiv:2608.08287
8. The Rustonomicon: Send and Sync — doc.rust-lang.org
9. RustBelt: Securing the Foundations of the Rust Programming Language — POPL 2018 PDF
10. Focus on Your Algorithm: NVIDIA CUDA Tile Handles the Hardware — NVIDIA Developer Blog