用Rust写CUDA内核?英伟达盖章了,真正的戏在条款里

源自66位全网作者

15:06

10月中旬,B站一条《老外炸了!英伟达官宣:GPU内核原生支持Rust直编译》的搬运视频开始刷屏。而更早聊 CUDA-Oxide 的视频区里,「CUDA 算子方向的又要失业了」这条评论拿了54个赞、排在第一。先给结论:官宣不在10月,这次焦虑有一半是错位的。哔哩哔哩哔哩哔哩

真实时间线

  • 5月,英伟达实验室开源实验性项目 CUDA-Oxide 0.1 编译器,让开发者用 Rust 直接开发 NVIDIA GPU 的 CUDA 内核,没有 DSL、没有 FFI、没有源到源转换。IT之家微博

  • 9月24日,英伟达官方专栏发布《介绍 CUDA Rust:用于编写 GPU 内核的两条轨道》,cuda-oxide 与 cutile-rs 两条原生路径就此官宣。知乎微博

  • 10月15日的中文搬运视频,只是二次传播,不是官宣当天。

用Rust写CUDA内核?英伟达盖章了,真正的戏在条款里

两条轨道是什么

  • 一条是底层 SIMT 轨道 cuda-oxide:函数标上 `#[kernel]`,经 MIR→Pliron IR→LLVM IR 直接编译出 PTX,但门槛不低——要 Linux、计算能力8.0起、CUDA 12.x,外加锁版本的 nightly Rust,英伟达自己都吐槽「这正是我们不想再向您提出的问题」。哔哩哔哩

  • 另一条是高层 Tile 轨道 cutile-rs:`#[cutile::module]` 宏把内核嵌进主机二进制,运行时靠 CUDA Tile IR JIT 编译,只要计算能力8.0起、CUDA 13.3、稳定版 Rust 1.89 和 Linux,不用 nightly、不用自建 LLVM。而且它已经有人生产使用——HuggingFace 的 Grout 推理引擎和 mistral.rs 都接上了。知乎

用Rust写CUDA内核?英伟达盖章了,真正的戏在条款里

真正解决的是什么
不是性能,是 kernel 里最噩梦的场景:竞争和越界。`DisjointSlice` 让每线程独占可变数据,编译期直接掐死非法别名,借用冲突当场编译失败、撞出 error[E0502]。但注意:SIMT 写共享内存仍然要 unsafe,Tile 模式则是没有线程索引的另一种写法。知乎微博

用Rust写CUDA内核?英伟达盖章了,真正的戏在条款里

用Rust写CUDA内核?英伟达盖章了,真正的戏在条款里

三拨人,三种声音
失业派喊「算子方向要凉了」。看条款:两个项目官方都标注仍处于早期阶段、未生产就绪。存量 C/C++ 内核一行都不用动——迁移只是"新代码可选",不是"老代码必搬"。知乎

"有什么区别"派拿27个赞:「都写 cuda 了,用啥语言有什么区别」。更尖锐的质疑是:llm kernel 用处不大,上 tcgen05 还得写 inline asm——这条并非没道理,Rust 轨道当下的卖点是内存安全和工程可维护性,不是性能。哔哩哔哩

护城河派看得最乐:「不敢相信这玩意的 backend 竟然开源,nvcc 都不开源,自毁护城河?」。高赞回复自己拆了台:“CUDA-Oxide = rustc + NVPTX,没什么好看的。” ISA 仍攥在英伟达手里,开源的是编译器后端,换的是开发者,护城河没塌,只是多了扇新门。哔哩哔哩

所以要不要学 Rust?
英伟达确实在把 Rust 潮水往自家工具链里引:官文顺带披露 Nova Linux 驱动用 Rust 写、Dynamo 核心基于 Rust、NVTX 有 Rust 绑定。但也别赌全量迁移——curl 那场 Rust 重构实验,4年跑到95%最终弃坑,就是新鲜的反面教材。务实建议就一条:别慌着转岗,但新起的小内核项目,可以把 Tile 轨道放进候选清单试试。知乎

彩蛋收尾:官方那篇《两条轨道》文内一度是机翻腔——「使用 Rust 有两个音轨,与 CUDA 本身的两个音轨相匹配」,评论区当场开嘲。护城河还是护城河,但官方号的翻译先漏了风。知乎

内容由AI生成
0
扫一下,分享更方便,购买更轻松
0评论

当前文章无评论,是时候发表评论了
提示信息

取消
确认
评论举报

最新文章 热门文章