cuda-oxide:hello-constant 全流程拆解——导读

用最简单的 hello-constant example,把 cuda-oxide 从一条 cargo 命令到 GPU 输出 42 之间发生的所有事,拆成 11 步逐站讲透。这是开篇——讲清楚我们要拆什么、为什么挑这个例子、整条路线长什么样。

zT
zTgx
HelloCUDA 站长
2026/06/11 3 周前 935
cuda-oxiderustcompilercuda ✓ 已验证
📚 系列 cuda-oxide · 第 4 篇

一条命令背后到底发生了多少事?这个系列接下来 11 篇,会用 cuda-oxide 自带的 hello_constant example,把 Rust 源码 → MIR → dialect IR → LLVM IR → PTX → SASS → GPU 输出 42 的整条链路一站一站拆开讲。这一篇是导读——交代起点、动机、路线图。

入口:一条命令

我们的起点就一条命令:

RUST_LOG=info cargo oxide run hello_constant

这命令一跑,过几秒钟终端会打出:

=========================================
RUSTC-CODEGEN-CUDA: hello_constant
=========================================

Target arch: sm_61 (auto-detected from CUDA device 0)
...
[PHASE 1/9] cargo-oxide driver started
[PHASE 2/9] spawning `cargo run --release` with codegen-backend RUSTFLAGS
[PHASE 3/9] rustc_codegen_cuda::codegen_crate — device code detected
[PHASE 4/9] collector: walked call graph from kernels
[PHASE 5/9] mir-importer::run_pipeline — MIR → dialect-mir → LLVM IR → PTX
[PHASE 6/9] translating MIR function to dialect-mir
[PHASE 7/9] mir-lower::lower_mir_to_llvm — dialect-mir → dialect-llvm
[PHASE 8/9] PTX file written (LLVM IR → PTX via llc)
[PHASE 9/9 a] host binary started
[PHASE 9/9 b] initializing CUDA context (device 0)
[PHASE 9/9 c] loading PTX module from hello_constant.ptx
[PHASE 9/9 d] launching kernel
[PHASE 9/9 e] copying device buffer back to host
Output: 42

九个 PHASE,从无到有出现了一个能在 GPU 上跑的程序。每个 PHASE 后面都藏着相当多的工程——编译器后端插件机制、MIR 翻译、dialect lowering、PTX 生成、driver JIT。这个系列要做的就是把这些细节挖出来。

这个 example 在做什么

打开 crates/rustc-codegen-cuda/examples/hello_constant/src/main.rs ,核心就两段:

#[kernel]
pub unsafe fn hello_constant(out: *mut i32) {
    let xxx = thread::xxx();
    gpu_printf!("thread xxx: {}", xxx);
    unsafe { *out = 42 };
}

fn main() {
    let ctx = CudaContext::new(0).expect("...");
    let stream = ctx.default_stream();
    let out_dev = DeviceBuffer::<i32>::zeroed(&stream, 1).expect("...");

    let cuda_module = ctx.load_module_from_file("hello_constant.ptx").expect("...");
    let module = kernels::from_module(cuda_module).expect("...");

    unsafe {
        module.hello_constant(
            stream.as_ref(),
            LaunchConfig::for_num_elems(1),
            out_dev.cu_deviceptr() as *mut i32,
        )
    }.expect("...");

    let result = out_dev.to_host_vec(&stream).expect("...");
    println!("Output: {}", result[0]);
}

每行做什么暂时不展开,这里先建立最粗的画面:

  • #[kernel] 标的函数最终会在 GPU 上跑
  • main() 是 host 代码,负责连 GPU、加载 kernel、launch、拷回结果
  • *out = 42 是 kernel 写的”载荷”(其实毫无意义,就是为了证明 kernel 跑过了)

为什么挑 hello_constant 拆

cuda-oxide 仓库里有好几个 example,从最简单的 vecadd 到工业级 6500 行的 gemm_sol。我挑 hello_constant 是因为:

优点解释
入门门槛低kernel 体只有 3 行,看完不会迷失
覆盖五个核心环节host 代码、kernel 注解、内建 intrinsic(thread::xxx)、device printf、内存写
跑 256 个线程自然引出 warp / block 概念(虽然每个线程做的事都一样)
不依赖 shared memory / 同步把 GPU 协作模型这层难点暂时屏蔽,聚焦”编译器流程”本身

更复杂的 example(tiled GEMM 那类)以后再单独开系列拆。先掌握 hello_constant ,所有更复杂的例子就只是”加更多 dialect op 和优化 pass”——骨架完全一样。

我们要解决什么疑问

这个系列要回答的问题清单:

  • cargo oxide run 跟普通 cargo run 区别在哪?额外干了什么?
  • rustc 怎么”知道”要走我们的 backend?.so 是怎么被加载进来的?
  • Rust 源码 → MIR 之间发生了什么?我们 backend 拿到的”原料”长什么样?
  • 怎么挑出哪些函数要编译成 GPU 代码?
  • thread::xxx() 这个内建函数实际是怎么被识别 + 替换的?
  • 多层 IR(MIR → dialect-mir → dialect-llvm → LLVM IR → PTX)为什么要这么多层?每一层换来了什么?
  • .ptx 不是机器码,GPU 是怎么真正执行它的?
  • 256 个线程在 GPU 上是怎么调度的?

读完整个系列,这些问题应该都有答案。

11 步路线图

我把整条流程切成 11 站。每篇一站,每站只讲一个核心概念,留时间消化:

篇号主题关键概念
05鸟瞰整个管线两个进程、7 层 IR 翻译链
06cargo-oxide driver配置 + spawn 薄层、-Z codegen-backend
07rustc 前端5 段流水线 → MIR(CFG of BB)
08codegen_crate 入口dlopen + 委托模式
09collector + stable MIRBFS 调用图、rustc_public 解耦
10MIR → dialect-mir(statement)alloca-load-store 模型 + mem2reg
11Terminator 翻译 + intrinsic dispatchcall 是 terminator、FQDN 字符串 match
12mir-lower(DialectConversion)op_cast 分发、partial vs full conversion
13llvm-export + llc文本化 + 外部 NVPTX backend
14CUDA driver 加载 PTXcuModuleLoad 触发 JIT(PTX → SASS)
15GPU 执行 + 结果回传256 线程 / 8 warp、benign race

两条贯穿全文的具体例子:*out = 42(普通 store)和 let xxx = thread::xxx()(自定义 intrinsic)。每一站都会回到这两条,看它们各自变成什么。

预备知识

如果对以下任何一项不熟,建议看完简介再开始:

  • rustc 基本架构:lexer / parser / HIR / MIR / codegen 这几层各自做什么
  • LLVM IR:基本指令、basic block、SSA 概念
  • MLIR 概念:dialect、op、interface(pliron 跟 MLIR 同构)
  • CUDA 基础:thread / warp / block / grid、kernel 是什么、__global__ 跟 __device__ 区别
  • PTX:NVIDIA 的虚拟 ISA,介于 LLVM IR 和真实 GPU 机器码 SASS 之间

不必精通——只要知道”它大致是什么”就够了。每篇会按需补充细节。

实战配套

写完每篇,推荐自己跑一遍 cuda-oxide,对照 dump 看实际数据。一条命令开齐所有 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 全套。对照博客一步步看,比单看文章效果好十倍。

一句话总结

这个系列接下来 11 篇,会从一条 cargo oxide run hello_constant 命令出发,把 Rust 源码到 GPU 输出 42 的整条编译 + 执行链路一站一站拆开,每站一个核心概念。hello_constant 的体型小到能扛得住这种密度,又五脏俱全足以覆盖 cuda-oxide 整套架构。下一篇:鸟瞰整个管线。

系列上一篇: cuda-oxide:用组合代替继承——一个真实例子

评论区
评论功能即将上线, 敬请期待。