写 GPU kernel 一直是 C++ 的地盘。cuda-oxide 把这件事换成 Rust——给 rustc 装一个特殊的代码生成后端,让
#[kernel]标记的 Rust 函数直接编译成 PTX,host 部分继续走正常 LLVM 路径。本系列拆它的内部架构。
1. cuda-oxide 是什么
一句话:一个 rustc 编译器后端,把带 #[kernel] 标记的 Rust 函数编译成 NVIDIA PTX。
具体一点:
#[kernel]
pub unsafe fn add_one(out: *mut i32) {
*out = 42; // ← 这函数最终在 GPU 上跑
}
fn main() {
// host 代码,在 CPU 上跑
let ctx = CudaContext::new(0).unwrap();
let buf = DeviceBuffer::<i32>::zeroed(&stream, 1).unwrap();
module.add_one(&stream, config, buf.ptr());
// ...
}同一个 .rs 文件里,kernel 和 host 代码共存:
#[kernel]函数走 cuda-oxide 路径 → PTX → CUDA driver JIT → 在 GPU 上跑- 其它代码走 rustc 自带 LLVM 后端 → x86_64 二进制 → 在 CPU 上跑
NVlabs 开源,定位是研究 / 教学级别的 toolchain,不追求跟 CUDA C++ 完全对等的性能,但够你拿来玩 PTX、写自己的 CUDA kernel、研究 GPU 编译器。
仓库:https://github.com/NVlabs/cuda-oxide
2. 为什么要用 Rust 写 CUDA
C++ 写 GPU kernel 已经存在三十年,Rust 进场的理由不能光靠”我喜欢 Rust”。几个有说服力的角度:
| 角度 | C++ kernel | Rust kernel |
|---|---|---|
| 类型安全 | 模板地狱 + 隐式转换 | trait + 显式类型 |
| 内存安全 | 手动指针,容易越界 | 借用检查 + unsafe 显式标 |
| 工具链 | nvcc + CMake | cargo + rustc |
| host/device 共享代码 | 头文件 + __host__ __device__ | crate + cfg |
| 错误处理 | error code 或 exception | Result + ? |
| 生态 | cuBLAS / cuDNN(C++) | cublas-rs、cudarc 等(刚起步) |
实话:Rust GPU 生态还远远没法跟 CUDA C++ 比。但作为研究方向,它打开了几个新可能:
- AI 编译器(Triton、IREE)开始用 Rust 写中间层,GPU 后端越来越重要
- 写 cuBLAS-rs / cuDNN-rs 这种纯 Rust 高性能库,需要直接生成 PTX 的能力
- rustc 的类型系统 + 借用检查 + zero-cost abstraction 在 GPU kernel 上是有意义的
cuda-oxide 不是终点,但它是Rust 直接编译到 GPU 这件事可行的证据。
3. cuda-oxide 的设计哲学
读懂 cuda-oxide 之前先建立三个心智模型——后面所有细节都是这三点的展开。
3.1 不重写 rustc,只插一个后端
rustc 有几十万行代码处理 Rust 语法、类型推断、借用检查、MIR 生成。cuda-oxide 完全不碰这些——它通过 -Z codegen-backend 这个 nightly flag 给 rustc 挂一个动态库后端,只在 codegen 阶段介入。
(-Z codegen-backend 的细节见 compiler 系列第二篇。)
3.2 Composition over Inheritance:host 路径委托给 LLVM 后端
我们的 backend 不替换整个 LLVM 后端——它包装一个 LlvmCodegenBackend 实例,绝大多数 trait 方法直接转发,只在 codegen_crate 这一处分叉:有 #[kernel] 走 cuda-oxide 管线,host 部分继续委托 LLVM 编。
这是 host + device “双轨制”的关键。详见 03 用组合代替继承。
3.3 Progressive Lowering:7 层 IR 一站一站降
device 路径不是 MIR 直接喷 PTX,而是经过 7 层 IR:
MIR → dialect-mir + dialect-nvvm → dialect-llvm → LLVM IR → PTX每一层比上一层更接近硬件、表达力更弱、约束更明确。这种”渐进式下降”是 MLIR / pliron 设计哲学的核心,也是 AI 编译器普遍采用的架构。
4. 本系列要讲什么
整个系列围绕两件事展开:
4.1 hello-constant 全流程拆解(主线)
从一条命令 cargo oxide run hello_constant 开始,把整条编译 + 执行链路切成 11 步,逐站讲透:
[04] 导读
[05] Step 01:鸟瞰整个管线
[06] Step 02:cargo-oxide driver 进程做了什么
[07] Step 03:rustc 前端把源码变成 MIR
[..] Step 04..11:codegen_crate / collector / mir-importer / mir-lower / llvm-export / llc / 运行时读完这条主线,你能从源码的 *out = 42 追到 GPU 上真正执行那条 SASS 指令。
4.2 单独主题(穿插在主线之间)
不属于”hello-constant 11 步”的小专题,但能加深对 cuda-oxide 的理解:
- [02] 确认 codegen backend 符号导出(用 nm 验证
.so) - [03] 用组合代替继承——一个真实例子
- [08] rustc backend dylib 是怎么编出来的
- (后续会有 WMMA wrapper、PTX 学习、tiled-gemm 拆解等)
编号规则:不管是主线”拆解”还是单独主题,统一按发布顺序往下排。读者从标题就能区分:带”hello-constant 拆解 XX”的是主线,其它的是单独主题。
5. 怎么跟这个系列
5.1 推荐顺序
- 先看本篇(开篇),理解 cuda-oxide 的定位
- 看 02 和 03 建立”backend 是动态库 + composition”的基本图景
- 从 04 导读 进入 hello-constant 主线,按 05、06、07 依次走
5.2 配套实战
每篇博客都建议自己跑一遍 cuda-oxide 对照看。一条命令开齐所有 dump:
RUST_LOG=rustc_codegen_cuda=debug,mir_importer=debug,mir_lower=debug,info \
CUDA_OXIDE_VERBOSE=1 \
CUDA_OXIDE_DUMP_MIR=1 \
CUDA_OXIDE_DUMP_LLVM=1 \
CUDA_OXIDE_SHOW_RUSTC_MIR=1 \
cargo oxide run hello_constant 2>&1 | tee /tmp/run.log输出包括 rustc MIR、dialect-mir(各阶段)、dialect-llvm、.ll、.ptx,对照博客一步步看,比单看文章效果好十倍。
5.3 预备知识
不必精通,知道大致是什么就够,每篇会按需展开:
- rustc 基本架构(MIR / codegen 阶段)
- LLVM IR(基本指令、basic block、SSA)
- MLIR 概念(dialect、op、interface,pliron 跟它同构)
- CUDA 基础(thread / warp / block / grid)
- PTX(NVIDIA 虚拟 ISA,介于 LLVM IR 和 SASS 之间)
6. 一句话总结
cuda-oxide 给 rustc 装了一个动态库后端,让带
#[kernel]的 Rust 函数直接编译成 PTX 跑在 GPU 上。本系列围绕一个最简单的 hello-constant example,把从源码到 GPU 输出 42 的整条编译 + 执行链路拆成 11 步逐站讲透,中间穿插单独主题——读完你能从 Rust 源码追到任意一条 PTX 指令的来源。
下一篇:确认 codegen backend 符号导出——用 nm 验证 backend 动态库的入口符号是否正确导出,迈出 cuda-oxide 工程入门的第一步。