CUDA Rust 两条路、Triton 两个补丁：从可运行样例到 MX 布局验证 (2026)

这期值得跟进的更新，集中在“能不能更安全地写 kernel”和“编译器能不能把边界条件看对”两件事上。CUDA Rust已经给出 SIMT 与 Tile 两条可运行路线；Triton 则补上了异步拷贝屏障追踪和 MX scale 布局的符号维度处理。三条材料都没有给出新的端到端吞吐提升，所以今天适合做小规模验证，不适合直接改线上性能结论。

先看结论

想在 Rust 里写 GPU kernel：先试 cutile-rs 的 Tile 路线；它使用 stable Rust 1.89+ 和 CUDA 13.3。需要显式线程、共享内存或更细的 launch 控制，再看 cuda-oxide 的 SIMT 路线，但它仍需要固定 nightly 和 LLVM。
正在用 Triton ConSan 检查异步 kernel：可以跟进 async-copy 与 mbarrier 的新追踪逻辑。提交说明给出了回归用例结果，但没有运行时加速数字。
正在处理 Blackwell、Hopper 或 CDNA4 的 MX scale：可以先验证符号 shape 查询。这个改动修的是布局计算的正确性和可表达性，不是已经测出的推理加速。

CUDA 与 GPU 编程

NVIDIA CUDA Rust：SIMT 和 Tile 各有一条原生路线

NVIDIA 在 9 月 8 日发布 CUDA Rust 介绍，核心变化不是“用 Rust 包一层 CUDA C++”，而是让 GPU kernel 本身直接从 Rust 编译到 PTX。它提供两条路线：cuda-oxide 走熟悉的 SIMT（每个线程执行什么），cutile-rs 走 Tile（每个数据块执行什么，由 Tile IR 编译器决定）。1

适用场景。 如果现有 kernel 需要精确控制线程索引、launch contract 或共享内存，SIMT 更接近 CUDA C++；如果 kernel 主要是张量块运算，希望编译器接管线程映射和内存布局，Tile 更适合作为第一条验证路径。NVIDIA 的原文也建议优先考虑 Tile，需要控制时再下沉到 SIMT。1

核心改动。 cuda-oxide 是定制的 rustc codegen backend，经 Rust MIR、Pliron 和 LLVM 生成 PTX；它要求 Linux、计算能力 8.0 及以上 GPU、CUDA 12.x 及以上、clang/libclang，以及固定 nightly 工具链。cutile-rs 把 kernel AST 嵌入 host binary，再通过 CUDA Tile IR JIT 编译；它要求 Linux、计算能力 8.0 及以上 GPU、CUDA 13.3 和 stable Rust 1.89+，不要求自带 LLVM。1

收益与边界。 两个示例都只验证了 1,024 个元素的向量加法结果，文章没有报告吞吐、延迟或与 CUDA C++ 的 benchmark。当前价值在于类型和所有权约束：cuda-oxide 用 DisjointSlice 与 launch contract 约束每线程写入范围，Tile 路线用 tensor partitioning 表达互斥写入。原文明确把 cuda-oxide 标为 early alpha，两个项目都还不是 production-ready。1

第一步验证。 先用你的目标 GPU 跑一次 cutile-rs 的 hello-world 或 vector-add，再记录首次 JIT 时间、后续缓存命中时间和结果校验；只有当依赖版本、驱动和 kernel 语义都稳定后，才值得把同一算子与现有 CUDA/Triton 实现做端到端对比。不要把示例的 PASSED 当成性能结论。

AI 编译器与 kernel 工具链

Triton ConSan：把 async-copy 的完成状态接到 mbarrier 上

Triton 在 9 月 9 日凌晨提交了 ConSan（并发 sanitizer）改动，主题是“通过 mbarrier 追踪 async-copy 完成”。改动针对 cp.async 写入在 cp.async.mbarrier.arrive 与 mbarrier.wait 之间仍处于 pending 的情况，把每个 barrier phase 关联的 copy 纳入已有 write-tracking，并在混用 wait_group 与 barrier 完成时保留 copy issuer，再把可见性转交给真正获取数据的线程/CTA。2 3

适用场景。 这条改动与使用 cp.async、shared memory、mbarrier，或把 TMA/copy completion 混在同一 kernel 里的 ConSan 检查直接相关。提交说明显示，在 NVIDIA GB200、ptxas 13.3.69 override 的回归中，55 个 GPU case 通过、6 个跳过、492 个取消选择，执行时间为 60.74 秒；NVIDIA 与 AMD 的 ConSan lit 回归文件也通过。这个 60.74 秒是测试执行时间，不是 kernel 加速。2

第一步验证。 如果你在用这类异步路径，先只跑提交说明列出的 async_copy、barrier invalidation 和 TMA multicast 相关 ConSan case，再检查报告是否仍把合法的 pending copy 判成 async_copy_global_to_shared 访问错误。当前来源没有给出吞吐或延迟变化，验证目标应是误报/漏报，而不是速度。

Triton MX scale：让符号维度留在整数域里

Triton 同一时间段的另一条 kernel 提交，修正 Blackwell、Hopper 和 CDNA4 上 MX scale storage shape 的符号整数计算。Blackwell 与 Hopper 的布局在加入对齐量前先处理 alignment - 1；CDNA4 则用整数向上取整替代浮点 ceil，因此即使 m、n 仍是符号维度，也能保持整数表达。提交新增的测试覆盖 BlackwellMXScaleLayout、BlackwellActMXScaleLayout、HopperMXScaleLayout 和 CDNA4MXScaleLayout，变更统计为 58 个测试项和 47 个非测试文件改动。3 4

适用场景。 当 kernel 或上层图编译器需要在逻辑维度尚未具体化时查询 MX scale 的 packed storage shape，这个改动能减少“先把符号量转成浮点”或手写架构分支的需要。它覆盖的是布局查询和测试可表达性；提交没有提供端到端推理 benchmark，也没有证明任何模型吞吐提升。

第一步验证。 在目标架构上用符号 m, n 调用公开 layout API，分别核对 Blackwell、Hopper、CDNA4 的 storage shape，再用具体整数维度做一次 materialized shape 对照。若你的调用只使用静态 shape，优先级较低；若模型存在 ragged batch、动态序列长度或 scale packing，优先检查这条路径是否仍触发类型转换错误。

今天怎么验证

Rust kernel：先选一条路线跑通向量加法，保存编译器、CUDA、驱动、GPU 型号和首次/后续运行时间；结果只作为 smoke test。
ConSan：对使用 cp.async + mbarrier 的现有 kernel 跑 targeted regression，确认报告从“合法 pending copy”与“真实越界/错误同步”之间分开。
MX scale：把动态维度和具体维度各跑一遍，检查 packed shape 与后续 load/store layout 是否一致；通过后再做模型级 benchmark。

今天的证据足以决定“哪些值得验证”，还不足以替读者决定“哪条一定更快”。

References

1
Introducing CUDA Rust: Two Tracks for Writing GPU Kernels
developer.nvidia.com
2
Triton commit：Track async-copy completion through mbarriers
github.com
3
Triton main 提交时间线
github.com
4
Triton commit：Keep scale layout shape calculations in integer arithmetic
github.com

CUDA Rust 两条路、Triton 两个补丁：从可运行样例到 MX 布局验证