系列前 16 篇按时间顺序拆完了 cuda-oxide 整套流程。这一篇不讲流程,直接把每个阶段的真实输出并排放出来。同一段 3 行 Rust 代码,在 7 层 IR 里长什么样——这是把抽象的”渐进式 lowering”变成具体感受的最快方式。

0. 比较哪 7 层

#	阶段	形态	由谁产生
1	Rust 源码	文本	用户写
2	rustc stable MIR	文本(单态化后)	rustc 前端
3	dialect-mir(刚翻译完,pre-verify)	pliron IR	mir-importer
4	dialect-mir(mem2reg 之后)	pliron IR(简化版)	pliron mem2reg pass
5	dialect-llvm	pliron IR	mir-lower
6	LLVM IR(.ll 文本)	textual LLVM IR	llvm-export
7	PTX(.ptx 文本)	NVPTX 汇编	llc(NVPTX backend)

实际跑 cuda-oxide 时这些产物全都能在 stderr / 文件里抓出来,本文所有片段都取自一次真实运行的输出。

0.1 “pre-verify” 和 “after mem2reg” 是什么意思

后面会反复出现两个标签 —— 提前说清楚:

mir-importer 内部的 pipeline:
        │
        ▼  translator(Step 6/7 干的事)
   dialect-mir 草稿     ← "pre-verify" 这个时刻
        │
        ▼  verify_operation 检查每个 op 是否合法
        │  (如果挂了,翻译过程出 bug;通过就继续)
        ▼
   verified dialect-mir
        │
        ▼  mem2reg 优化 pass
        │
        ▼
   优化后的 dialect-mir   ← "after mem2reg" 这个时刻
        │
        ▼  mir-lower(Step 8 干的事)
   dialect-llvm

pre-verify = translator 刚把 stable MIR 翻译成 dialect-mir 的”草稿”状态,verify 还没跑。dump 在这个时刻是为了:如果 verify 挂了,能让用户看到导致失败的 IR 长什么样。verify 通过后,这份草稿就直接进 mem2reg。

after mem2reg = mem2reg pass 跑完之后的状态,这才是接下来交给 mir-lower 的真正输入。mem2reg 把冗余的栈分配消掉(下面 §4 详解),所以这一版通常比 pre-verify 短很多。

Part 1:整个 hello_kernel 横跨 7 层

hello_kernel 是 example 里最简单的 kernel——Rust 源码 3 行,PTX 6 条指令。全函数从源码到 PTX 的演化能一眼看全。

1. Rust 源码

#[kernel]
pub unsafe fn hello_kernel(ins: i32, out: *mut i32) {
    *out = ins + 1;
}

3 行,函数体一句话。

2. rustc stable MIR

fn hello_kernel {
    let mut _0: ();
    let _1:  i32;
    let _2:  *mut i32;
    debug "ins" => _1;
    debug "out" => _2;
    bb0: {
        (*_2) = Add(copy _1, const 1_i32)
        return
    }
}

注意几个细节(这些后面每一层都会渐次”消失”):

现象	含义
`let _1: i32; let _2: *mut i32`	rustc 给每个 local 编号,显式声明类型
`debug "ins" => _1`	保留源码的变量名映射(只为 debug info,不影响 codegen)
`(*_2) = Add(copy _1, const 1_i32)`	Rvalue 是嵌套表达式(`Add(_, _)`)——比 SSA 形式紧凑
`bb0`、`return` 没分号	rustc MIR pretty-print 格式

整段就 1 个 basic block + 2 条 statement + 1 个 terminator。

3. dialect-mir(pre-verify,mir-importer 刚翻译完)

mir.func @hello_kernel: builtin.function <(builtin.integer si32, mir.ptr <builtin.integer si32,mutable:true,addrspace:0>)->()>
  [gpu_kernel: builtin.string "true"]
{
  ^block5v1(v40: builtin.integer si32, v41: mir.ptr <builtin.integer si32,mutable:true,addrspace:0>):
    v42 = mir.alloca () ... -> (mir.ptr <builtin.integer si32,mutable:true,addrspace:0>);
    v43 = mir.alloca () ... -> (mir.ptr <mir.ptr <builtin.integer si32,mutable:true,addrspace:0>,mutable:true,addrspace:0>);
    mir.store (v42, v40) ...;
    mir.store (v43, v41) ...;
    v44 = mir.load (v42) ... -> (builtin.integer si32);
    v45 = mir.constant () ... [value: builtin.integer <1: si32>] -> (builtin.integer si32);
    v46 = mir.add (v44, v45) ... -> (builtin.integer si32);
    v47 = mir.load (v43) ... -> (mir.ptr <builtin.integer si32,mutable:true,addrspace:0>);
    mir.store (v47, v46) ...;
    mir.return ()
}

新东西:

现象	含义
`^block5v1(v40, v41)`	block 用 `^name` 标号,参数显式声明
`mir.alloca` ×2	给两个参数各开一个栈槽——经典 alloca-load-store 模型
两条 `mir.store (slot, param)`	把入参拷进栈槽
`mir.add(v44, v45)` 不再是嵌套表达式	拆成 op,扁平化
`builtin.integer si32`	有符号性显式写出来(`si32` = signed 32-bit)
`mir.ptr <... mutable:true, addrspace:0>`	指针带可变性 + 地址空间
`[gpu_kernel: builtin.string "true"]`	kernel 标记,一路传到 PTX `.entry`

整段 10 个 op,膨胀了。是 mem2reg 还没跑的”草稿”形态。

4. dialect-mir(mem2reg 之后)

mir.func @hello_kernel: builtin.function <(builtin.integer si32, mir.ptr <builtin.integer si32,mutable:true,addrspace:0>)->()>
  [gpu_kernel: builtin.string "true"]
{
  ^block5v1(v40: builtin.integer si32, v41: mir.ptr <builtin.integer si32,mutable:true,addrspace:0>):
    v45 = mir.constant () ... [value: builtin.integer <1: si32>];
    v46 = mir.add (v40, v45) ...;
    mir.store (v41, v46) ...;
    mir.return ()
}

10 个 op → 4 个 op。这一节细看:mem2reg 干了什么、为什么干、对 hello_kernel 的具体影响。

4.1 mem2reg 一句话

看每个 mir.alloca,如果它的地址只被 mir.load 和 mir.store 直接用、没”逃逸”到别处,就把这个 alloca 整个消掉:所有 mir.store slot, v 删掉,所有 mir.load slot 替换成最近一次写进 slot 的 SSA valuev。

“地址逃逸”指的是 alloca 的指针被传给别的函数、被 cast 成其它类型、被存进别的指针——只要它不是被 load/store 直接消费,就不能 promote(因为编译器没法追踪它会被改成什么)。

4.2 hello_kernel 里哪些 alloca 是 promotable

pre-verify 版本里有两个 alloca:

v42 = mir.alloca   → ptr<si32, mut>             ← _1 (ins) 的栈槽
v43 = mir.alloca   → ptr<ptr<si32, mut>, mut>   ← _2 (out) 的栈槽

它们的所有 use:

alloca	唯一的 store	唯一的 load
`v42`	`mir.store (v42, v40)` ← 把参数 `v40` 拷进去	`v44 = mir.load (v42)`
`v43`	`mir.store (v43, v41)` ← 把参数 `v41` 拷进去	`v47 = mir.load (v43)`

每个 alloca 只被写一次,只被读一次,地址没传给任何其它东西——典型的 promotable。

4.3 mem2reg 怎么改写

对 v42:

找到唯一的 store:mir.store (v42, v40) ——这告诉 mem2reg “v42 这个 slot 持有的值是 v40”
找到唯一的 load:v44 = mir.load (v42) ——把所有用到 v44 的地方直接换成 v40
删掉 store、删掉 load、删掉 alloca

v43 同理,所有用到 v47 的地方换成 v41。

4.4 影响的具体语句

pre-verify 里有一行:

v46 = mir.add (v44, v45) ...
                ↑
                这是 mir.load (v42) 的结果

mem2reg 把 v44 换成 v40:

v46 = mir.add (v40, v45) ...
                ↑
                直接用参数本身

还有一行 store 用到 v47(原来是 mir.load (v43) 的结果):

mir.store (v47, v46) ...   →   mir.store (v41, v46) ...
              ↑                              ↑
              load v43 的结果              直接用参数 v41

4.5 消失的 6 条 op

被消的 op	为什么
`v42 = mir.alloca`	整个槽不要了
`v43 = mir.alloca`	整个槽不要了
`mir.store (v42, v40)`	没人需要”把 v40 写进 v42 槽”了
`mir.store (v43, v41)`	同上
`v44 = mir.load (v42)`	所有用 v44 的地方换成 v40
`v47 = mir.load (v43)`	所有用 v47 的地方换成 v41

剩下 4 条 op:constant 1 + add + 最后的 store *out = result + return。

4.6 为什么这件事这么重要

如果 mem2reg 不跑,后面 lower 到 PTX 时,这两个 alloca 会变成 PTX 的 .local 状态空间(外置 DRAM)。每个 ld.local / st.local 访问延迟比寄存器慢几十倍。

跑了 mem2reg,参数直接以寄存器形态流过整个函数 —— PTX 里就是干净的 ld.param(读参数)→ add(算)→ st.global(写)三步,没有任何 .local 出现。

这就是为什么 mem2reg 在 cuda-oxide 是”必经一步”而不是”可选优化”。

5. dialect-llvm(mir-lower 降级后)

llvm.func @hello_kernel: llvm.func <llvm.void (builtin.integer i32, llvm.ptr addrspace(0)) variadic = false>
  [gpu_kernel: builtin.string "true"]
{
  ^entry_block7v1(v84: builtin.integer i32, v85: llvm.ptr addrspace(0)):
    llvm.br ^block5v1(v84, v85)

  ^block5v1(v40: builtin.integer i32, v41: llvm.ptr addrspace(0)):
    v86 = llvm.constant <builtin.integer <1: i32>> : builtin.integer i32;
    v87 = llvm.add v40, v86 <{nsw=false,nuw=false}>: builtin.integer i32;
    llvm.store *v41 <- v87;
    llvm.return
}

发生了什么:

变化	含义
`si32` → `i32`	有符号性消失——LLVM IR 没有 signed/unsigned 概念,只有位宽
`mir.ptr<si32, mutable:true, addrspace:0>` → `llvm.ptr addrspace(0)`	指针类型擦除,只剩地址空间(LLVM 现在用 opaque pointer)
`mir.add(v40, v45)` → `llvm.add v40, v86 <{nsw=false, nuw=false}>`	有符号性挪进操作属性(`nsw` = no signed wrap,`nuw` = no unsigned wrap)
多了 `^entry_block7v1` 入口块 + `llvm.br`	mir-lower 在每个函数前插一个 entry block,转发参数到原 block

[gpu_kernel: "true"] attribute 仍然保留,接力到下一层。

6. LLVM IR(`.ll` 文本)

define ptx_kernel void @hello_kernel(i32 %v0, ptr %v1) {
entry:
  br label %bb0
bb0:
  %v2 = phi i32 [ %v0, %entry ]
  %v3 = phi ptr [ %v1, %entry ]
  %v4 = add i32 %v2, 1
  store i32 %v4, ptr %v3
  ret void
}

文本化之后:

变化	含义
`ptx_kernel`	calling convention 关键字,告诉 NVPTX backend 这是 GPU 入口
`%v0 %v1 %v2 ...`	LLVM SSA value 命名风格,`%` 前缀
`entry:` + `bb0:` 标签	LLVM IR 的 block 用 `label:` 形式
`phi i32 [ %v0, %entry ]`	入口转发被翻成 phi 节点——SSA 形式的”汇合点”
`add i32 %v2, 1`	LLVM 自动把 `mir.constant 1` 内联进 `add`
`ret void`	LLVM IR 的 return

注意:gpu_kernel: "true" attribute 在文本 IR 里变成 ptx_kernel 这个 calling convention(或者 !nvvm.annotations 元数据,看具体路径)。

7. PTX(`.ptx` 文本)

.visible .entry hello_kernel(
    .param .u32 hello_kernel_param_0,
    .param .u64 .ptr .align 1 hello_kernel_param_1
)
{
    .reg .b32 %r<3>;
    .reg .b64 %rd<3>;

    ld.param.b32 %r1, [hello_kernel_param_0];
    ld.param.b64 %rd2, [hello_kernel_param_1];
    cvta.to.global.u64 %rd1, %rd2;
    add.s32 %r2, %r1, 1;
    st.global.b32 [%rd1], %r2;
    ret;
}

完全不一样的形态:

变化	含义
`.visible .entry`	PTX 关键字:对外可见 + kernel 入口
`.param .u32 / .u64`	参数存放在 parameter state space——这是 PTX 的概念,LLVM IR 里看不到
`.reg .b32 %r<3>`	显式声明用了几个寄存器(`%r1` `%r2`,`<3>` 是上界)
`ld.param.b32 %r1, [...]`	把参数从 `.param` 空间 load 到寄存器
`cvta.to.global.u64`	通用地址 → global 地址空间 —— 这是 PTX 独有
`add.s32 %r2, %r1, 1`	有符号性又回来了,挂在 opcode 上(`.s32`)
`st.global.b32 [%rd1], %r2`	写到 global memory

8. 七层叠在一起:行数/op 数对比

层	函数体大小	关键变化
1 Rust 源码	1 行函数体	`*out = ins + 1;`
2 rustc MIR	1 个 bb + 2 条 statement	嵌套 Rvalue,显式 local 编号
3 dialect-mir(pre)	10 个 op	alloca-load-store 膨胀
4 dialect-mir(post-mem2reg)	4 个 op	alloca 消失
5 dialect-llvm	4 个 op + 入口 br	类型擦除,有符号性挪到操作属性
6 LLVM IR	5 行 SSA	文本 SSA + phi
7 PTX	6 条指令	寄存器、state space、显式 load/store

Rust 1 行 → MIR 2 statement → dialect-mir 4 op → PTX 6 条指令。信息一层层从”做什么”变成”怎么做”。

Part 2:单语句的逐层追踪

整函数看完了,放大看两条具体语句——它们代表两种典型路径。

9. 路径 A:`*out = 42`(普通 store)

源码 1 行:*out = 42;(在 hello_constant 里)

A.1 rustc stable MIR

bb2: {
    ...几个 StorageDead...
    (*_1) = const 42_i32
    return
}

一条 statement,Place((*_1))+ Rvalue(const 42_i32)。

A.2 dialect-mir(pre-verify)

v38 = mir.constant () ... [value: builtin.integer <42: si32>];
v39 = mir.load (v1) ... -> (mir.ptr <builtin.integer si32, mutable:true>);
mir.store (v39, v38) ...;

3 条 op:

v38 = 常量 42
v39 = 从 _1 栈槽 load 出 out 指针
store 把 42 写到指针指向的地方

A.3 dialect-mir(mem2reg 之后)

v38 = mir.constant () ... [value: builtin.integer <42: si32>];
mir.store (v0, v38) ...;

2 条 op——mir.load (v1) 消失了,因为 mem2reg 发现 v1 槽只在入口被参数 v0 写过一次,后面所有 load 直接换成 v0。

A.4 dialect-llvm

v83 = llvm.constant <builtin.integer <42: i32>> : builtin.integer i32;
llvm.store *v0 <- v83;

跟 dialect-mir 一致,只是类型 si32 → i32。

A.5 LLVM IR(`.ll`)

store i32 42, ptr %v1

就一条。LLVM 进一步把常量 42 直接内联到 store——v83 = constant 42 这个临时也省了。

A.6 PTX

st.global.b32 [%rd1], 42;

一条 PTX 指令。常量 42 又再次内联(从 LLVM IR store i32 42 直接映射)。

A.7 一条 statement 的演化总览

层	形态	op/指令数
Rust	`*out = 42`	1
MIR	`(*_1) = const 42_i32`	1 statement
dialect-mir(pre)	constant + load + store	3 ops
dialect-mir(post-mem2reg)	constant + store	2 ops
dialect-llvm	constant + store	2 ops
LLVM IR	`store i32 42, ptr %v1`	1
PTX	`st.global.b32 [%rd1], 42;`	1

形态像 U 形——MIR 翻译后先膨胀到 3 ops(因为 alloca-load-store 模型),mem2reg 把它压缩回 2 ops,llvm-export + llc 又进一步压缩到 1 条指令。

10. 路径 B:`let xxx = thread::xxx()`(intrinsic dispatch)

这条路径最有意思——桩函数体根本没翻译,直接被替换。

B.1 rustc stable MIR

bb0: {
    _2 = cuda_device::xxx() -> [return: bb1, unwind continue]
}

一个 Call terminator,它本身就是块尾。

B.2 dialect-mir(pre-verify)

v11 = nvvm.read_ptx_xxx () ...: <() -> (builtin.integer ui32)>;
mir.store (v2, v11) ...;
mir.goto () [^block3v1] ...;

3 条 op:

mir-importer 不翻译 xxx() 函数体,直接发 nvvm.read_ptx_xxx(一个 dialect-nvvm op)
写到 _2 槽
branch 到 return target(原 MIR 的 bb1)

B.3 dialect-mir(mem2reg 之后)

v11 = nvvm.read_ptx_xxx () ...;
mir.goto () [^block3v1] ...;

mir.store (v2, v11) 消失了——_2 这个槽在后续 block 只被 load 一次,mem2reg 直接把所有 mir.load (v2) 换成 v11。

B.4 dialect-llvm

v49 = llvm.inline_asm ...;
llvm.br ^block3v1();

nvvm.read_ptx_xxx 通过 MirToLlvmConversion::convert 翻译成 llvm.inline_asm op(因为用户当初写的就是 inline_asm_convergent("mov.u32 $0, 42;", "=r"))。

B.5 LLVM IR(`.ll`)

%v2 = call i32 asm sideeffect "mov.u32 $0, 42;", "=r"() #0

LLVM IR 的 inline asm 形式——sideeffect 是 LLVM 自动加的,告诉优化器有副作用别动。

B.6 PTX

// begin inline asm
mov.u32 %r1, 42;
// end inline asm

NVPTX backend 把 inline asm 字符串原封不动嵌进 PTX,前后加 // begin inline asm / // end inline asm 包裹。

B.7 一条 intrinsic 调用的演化总览

层	形态	op/指令数
Rust	`let xxx = thread::xxx()`	1
MIR	`_2 = cuda_device::xxx() -> [return: bb1, unwind continue]`	1 Call terminator
dialect-mir(pre)	NVVM op + store + goto	3 ops
dialect-mir(post-mem2reg)	NVVM op + goto	2 ops
dialect-llvm	inline_asm + br	2 ops
LLVM IR	`call i32 asm sideeffect ...`	1 行
PTX	`mov.u32 %r1, 42;`	1 条

关键观察:桩函数体(unreachable!())从来没翻译过——它在 MIR 里存在,但 try_dispatch_intrinsic 在 Step 7 那一步直接绕过去,把整个 Call 替换成 NVVM op。整条路径从来没碰过 xxx 函数的内部 MIR。

Part 3:几个维度的演化对照

把前面所有现象按维度归纳一下。

11. 操作数命名的演化

层	命名风格
Rust 源码	变量名(`ins`、`out`、`xxx`)
MIR	编号 `_0 _1 _2`(`debug "name" => _N` 保留名字映射)
dialect-mir	`v0 v1 v40 v41` ——pliron 全局唯一编号,顺序生成
dialect-llvm	同 `v40 v41 v86 v87` ——继承自 dialect-mir
LLVM IR	`%v0 %v1 %v2` ——加 `%` 前缀,函数内编号
PTX	`%r1 %r2 %rd1` ——按类型分前缀(`%r` = 32-bit,`%rd` = 64-bit)

12. 类型表达的演化

层	`*mut i32` 怎么写
Rust 源码	`*mut i32`
MIR	`*mut i32`
dialect-mir	`mir.ptr <builtin.integer si32, mutable:true, addrspace:0>`
dialect-llvm	`llvm.ptr addrspace(0)` ——指针类型擦除(opaque pointer)
LLVM IR	`ptr`
PTX	不再显式存在——参数声明用 `.u64 .ptr`,操作用 `.global` state space

有符号性(si32 vs i32):dialect-mir 之前显式标在类型上,dialect-llvm 起挪进操作属性(nsw/nuw),PTX 又回到 opcode(add.s32 等)。

13. 内存模型的演化

层	写一个值
Rust	`*out = 42`
MIR	`(*_1) = const 42_i32` ——一条 statement
dialect-mir(pre-verify)	constant + load(load 槽里的指针)+ store 三条 op
dialect-mir(post-mem2reg)	constant + store ——alloca 消失
LLVM IR	`store i32 42, ptr %v1` ——常量内联到 store
PTX	`st.global.b32 [%rd1], 42;` ——显式 global state space

地址空间这个概念,在 Rust / MIR 阶段不存在,dialect-llvm 开始有 addrspace(N),PTX 用 .global / .shared / .local / .param 等关键字显式标。

14. Intrinsic 表达的演化

层	`thread::xxx()` 怎么表示
Rust	`let xxx = thread::xxx();`
MIR	`_2 = cuda_device::xxx() -> [return: bb1, unwind continue]`(普通 Call)
dialect-mir	`nvvm.read_ptx_xxx ()`(专用 NVVM op,不走普通 call)
dialect-llvm	`llvm.inline_asm`
LLVM IR	`call i32 asm sideeffect "mov.u32 $0, 42;", "=r"()`
PTX	`mov.u32 %r1, 42;`(inline asm 原文嵌进去)

整条路径上,桩函数(unreachable!())的 MIR 从来没被翻译——它在 MIR 里只是一个名字,被 try_dispatch_intrinsic 直接绕过。

Part 4:总览

15. 信息怎么逐层丢失

阶段过渡	丢了什么
Rust → MIR	闭包、宏、`for/while` 语法糖、trait 解析
MIR → dialect-mir	`debug` 名字映射、Rvalue 嵌套结构(被拆扁)
dialect-mir → dialect-llvm	有符号性(挪到操作属性)、指针类型(变 opaque)
dialect-llvm → LLVM IR	`gpu_kernel` 属性(变成 calling convention)
LLVM IR → PTX	SSA 形式(变成寄存器复用)、phi 节点(变成 block 入参)

16. 形态怎么逐层接近硬件

阶段过渡	增了什么
Rust → MIR	basic block + CFG、单态化、显式 local 编号
MIR → dialect-mir	alloca-load-store 模型、显式 op 拆扁
dialect-mir → dialect-llvm	`addrspace(N)`、入口 entry block
dialect-llvm → LLVM IR	phi 节点(SSA 跨 block 汇合)、`ptx_kernel` calling convention
LLVM IR → PTX	state space(`.global` `.param` `.local`)、寄存器命名(`%r %rd`)、地址空间转换(`cvta`)、参数槽布局

17. 一句话总结

同一段 3 行 Rust 代码,从 rustc MIR 的 1 个 bb 2 条 statement,膨胀到 dialect-mir(pre-verify)的 10 个 op,mem2reg 之后压回 4 个 op,dialect-llvm 还是 4 个 op(但有符号性挪到 nsw),LLVM IR 变成 5 行 SSA + phi,最后落到 PTX 是 6 条指令。每一层都在用自己的语言表达同一件事——*out = ins + 1——但语言一步步从”做什么”(Rust 语义)磨成”怎么做”(state space 寄存器存取)。U 形演化(膨胀 → 收缩 → 平铺)是这套渐进式 lowering 设计的形状。看清楚这张对照表,前面 16 篇拆下来的所有抽象概念都能落到具体的字符串上。

系列上一篇: cuda-oxide:hello-constant 拆解 11(完结)——GPU 执行 + 结果回传

评论区

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

cuda-oxide:同一段代码在七层 IR 里长什么样——hello_constant 全程对比