上一篇(找 #[kernel] + BFS + stable MIR)把 stable
Instance喂给mir_importer::run_pipeline。这一篇拆 pipeline 内部:逐函数翻译 stable MIR 成 pliron 的dialect-mir表示。重点说清楚alloca-load-store这个核心设计模型,以及为什么后面mem2reg是必须而不是可选优化。
1. Pipeline 启动:创建空 pliron module
pub fn run_pipeline(functions, device_externs, config) -> Result<...> {
let mut ctx = Context::new(); // ① pliron 全局上下文
crate::translator::register_dialects(&mut ctx); // ② 注册 dialect
let module = pliron::builtin::ops::ModuleOp::new(&mut ctx, module_name);
let module_op_ptr = module.get_operation(); // ③ 空 module op
for func in functions { // ④ 逐函数翻译
let body = func.instance.body()?; // 单态化 MIR
let func_op_ptr = translator::body::translate_body(
&mut ctx, &body, &func.instance,
func.is_kernel, Some(&func.export_name), ...
)?;
verify_operation(&ctx, func_op_ptr, &func.export_name)?;
append_to_module(&ctx, module_op_ptr, func_op_ptr);
}
verify_operation(&ctx, module_op_ptr, "module")?;
mem2reg(module_op_ptr, &mut ctx)?; // ⑤ 优化
lower_to_llvm(&mut ctx, module_op_ptr)?; // ⑥ 进入下一站
// ...
}三个关键概念:
| 概念 | 角色 |
|---|---|
Context | pliron 的”心脏”。所有 Operation / Value / Type / Attribute 都在它里面存储,外界只持有 Ptr<T> 轻量句柄(类似 LLVM 的 LLVMContext) |
ModuleOp | 顶层容器,对应一个翻译单元,里面装多个 FuncOp |
| 三个 dialect | dialect-mir(保留 Rust 语义的中间层)、dialect-nvvm(GPU intrinsic)、dialect-llvm(下降目标) |
2. translate_body:翻译单个函数
输入是 stable MIR Body(基本块 + locals + signature),输出是一个 pliron FuncOp。翻译流程四步:
fn translate_body(ctx, body, instance, is_kernel, export_name, ...) -> FuncOp {
// ① 翻译函数签名(每个参数类型从 Ty → pliron Type)
let func_type = translate_signature(ctx, body);
// ② 创建空 FuncOp + entry basic block
let func_op = FuncOp::new(ctx, export_name, func_type, is_kernel);
let entry_block = func_op.entry_block();
// ③ 为每个 MIR local 在 entry 创建 alloca(栈槽)
let mut value_map = ValueMap::new();
for local in body.locals() {
let slot = mir::AllocaOp::new(ctx, local.ty); // mir.alloca <ty>
value_map.register_local(local.id, slot);
}
// ④ 逐个 basic block 翻译
for bb in body.blocks() {
let block_ptr = func_op.get_or_create_block(bb.id);
translate_block(ctx, body, bb, &mut value_map, block_ptr, ...);
}
func_op
}第 ③ 步是这个 pipeline 设计的核心选择,下面单独讲。
3. 关键设计:alloca-load-store 模型
每个 MIR local _1、_2、_3 都被翻译成一个栈槽(mir.alloca),不是直接的 SSA value。
为什么?因为 MIR 允许跨 block 反复写同一个 local:
bb0:
_2 = const 1
goto bb1
bb1:
_2 = const 2 ← _2 在不同 block 被重新赋值
goto bb2
bb2:
return _2 ← 这里 _2 到底是 1 还是 2?要做成纯 SSA 必须插 φ (phi) 节点,跨 block 数据流分析很麻烦。直接用 alloca + store + load,把 _2 当成”一个内存位置”,读时 load、写时 store——简单粗暴,翻译时不用考虑 phi。
代价:产生一堆冗余的 alloca / load / store。这就是 mem2reg 优化要解决的(第 7 节)。
4. 真实 dialect-mir 长什么样
跑一遍 hello_constant,run.log 里 hello_constant 函数翻译完的 dialect-mir 头部(简化注释):
mir.func @hello_constant: builtin.function <(mir.ptr<si32,mut,addrspace:0>)->()>
[gpu_kernel: true]
{
^block2v1(v0: mir.ptr<si32, ...>): ← entry block,v0 是 *mut i32 参数
v1 = mir.alloca () : ... -> *(*mut i32) ← _1 槽(参数 out)
v2 = mir.alloca () : ... -> *(u32) ← _2 槽(xxx)
v3 = mir.alloca () : ... -> *(i32) ← _3 槽(printf 返回值)
v4 = mir.alloca () : ... -> *(__GpuPrintfArgs) ← _4 槽(printf args)
v5 = mir.alloca () : ... -> *(i64) ← _5..._10 槽,省略
v6 = mir.alloca () ...
v7 = mir.alloca () ...
v8 = mir.alloca () ...
v9 = mir.alloca () ...
v10 = mir.alloca () ...
mir.store (v1, v0) ← *_1_slot = 参数 v0
v11 = nvvm.read_ptx_xxx () ← thread::xxx() 直接发 NVVM op
mir.store (v2, v11) ← *_2_slot = v11
mir.goto ^block3v1 ← 进入 bb1
^block3v1():
... (printf 参数构造,十几条 storage_live / cast / store / construct_struct) ...
v37 = nvvm.vprintf (v35, v36) ← 调 vprintf
mir.store (v3, v37)
mir.goto ^block4v1 ← 进入 bb2
^block4v1():
... (storage_dead 清理几个 local) ...
v38 = mir.constant 42_si32 ← 常量 42
v39 = mir.load (v1) ← 读 out 指针
mir.store (v39, v38) ← *out = 42 ★
mir.return
}几个观察:
| 观察 | 说明 |
|---|---|
| entry block 头部全是 alloca | 10 个 local 对应 10 条 mir.alloca,占了整整 10 行 |
mir.store (v1, v0) | 进 entry 第一件事:把函数参数 v0 copy 进 v1 这个栈槽 |
v11 = nvvm.read_ptx_xxx () | thread::xxx() 被 importer 直接换成 dialect-nvvm op,桩函数体没翻译(下一篇拆 terminator + intrinsic dispatch 时讲) |
mir.goto | Call terminator 翻译完的结果(NVVM op + 跳到 return target block) |
| 3 个 block 对应原 MIR 的 3 个 bb | 每次 Call 切一刀,所以源码 3 行变 3 block |
*out = 42 在 ^block4v1 末尾 | 由 constant + load + store 三条 op 实现 |
5. 拆 *out = 42 一条 statement 的翻译过程
源码:
*out = 42;stable MIR statement(简化):
Statement {
kind: Assign(
place: Place {
local: _1, // out 参数
projection: [Deref], // *_1
},
rvalue: Use(Operand::Constant(42_i32)),
),
}translator 处理 Assign 分支:
match &stmt.kind {
mir::StatementKind::Assign(place, rvalue) => {
// ① 翻译 rvalue 拿到要赋的值
let (rvalue_op_opt, result_value, last_inserted) =
rvalue::translate_rvalue(ctx, body, rvalue, value_map, ...)?;
// ② 根据 place 的 projection 决定怎么写
if place.projection.is_empty() {
// _2 = ... (普通 local 赋值)
value_map.store_local(ctx, place.local, result_value, ...)
} else if matches!(place.projection[0], ProjectionElem::Deref) {
// *_1 = ... ← 我们的情况
// ...先 load _1 拿到指针,再 store value 到指针指向的位置
}
}
}5.1 翻译 rvalue Use(Constant(42))
Rvalue::Use(Operand::Constant(42_i32)) 不发新的 op,只创建一个字面量 SSA Value。run.log 里对应:
v38 = mir.constant () [] [value: builtin.integer <42: si32>]mir.constant op 把 42 装进一个 Value,返回 result_value = v38。
5.2 翻译 place *_1
ProjectionElem::Deref 触发”先 load 指针,再 store value”两步:
v39 = mir.load (v1) ← _1 槽里存的是 out 指针,load 出来
mir.store (v39, v38) ← *v39 = v38,即 *out = 425.3 完整的三条 op
合起来 *out = 42 实际生成:
v38 = mir.constant 42_si32 ← rvalue 求值
v39 = mir.load (v1) ← 读 _1 栈槽(拿 out 指针)
mir.store (v39, v38) ← 写 *out = 42注意 entry 里还有一条 mir.alloca v1,所以总共四步:一个 alloca(entry 里建)+ 三条具体 op(写值时发)。
6. Statement 和”块尾控制流”的边界
打开 run.log 看 hello_constant 的 stable MIR,每个 bb 内部都是这样:
bb2: {
StorageDead(_8)
StorageDead(_6)
StorageDead(_4)
StorageDead(_3)
(*_1) = const 42_i32
return ← 块的最后一行,决定下一步去哪
}bb 里前面是 statement,最后一行是控制流——rustc 内部把后者叫 Terminator(mir::Terminator),但 dump 里没有这个标签,就是最后那一行的形态告诉你它是什么。
hello_constant 在 run.log 里实际出现的几种块尾:
bb0: {
_2 = cuda_device::xxx() -> [return: bb1, unwind continue]
}
bb1: {
...
_3 = cuda_device::debug::__gpu_vprintf(move _6, move _8) -> [return: bb2, unwind continue]
}
bb2: {
...
(*_1) = const 42_i32 ← 这是 statement,不是块尾
return ← 这才是块尾
}几种常见块尾形态:
| 真实 dump 形态 | rustc 内部类型名 | 含义 |
|---|---|---|
_N = some_fn(...) -> [return: bbM, unwind continue] | Call | 调函数,正常返回到 bbM |
goto -> bbM | Goto | 无条件跳到 bbM |
switchInt(_N) -> [0: bb1, otherwise: bb2] | SwitchInt | if / match 编译后的多路分支 |
return | Return | 函数返回 |
unreachable | Unreachable | 不可达 |
注意三个真实细节:
- 块尾不带分号,跟 statement 一样,整个 MIR dump 都不带
; unwind continue中间没有冒号——continue是UnwindAction枚举值(token 形式),跟return: bb1的 label 形式不一样(*_1) = const 42_i32是 statement 不是块尾——别被它的格式迷惑
translate_block 先逐个翻译前面的 statement(就是这一篇讲的 mir.alloca / mir.load / mir.store 等),最后翻译块尾控制流——下一篇专门拆。
这一篇已经在 run.log 里见过的 statement-级 op:mir.alloca、mir.store、mir.load、mir.constant、mir.cast、mir.construct_struct、mir.construct_array、mir.ref、mir.storage_live / storage_dead。块尾的 mir.goto、mir.return 留给下一篇。
7. mem2reg 之后会变成什么
mem2reg 是一个经典 SSA 重建优化。它扫整个函数,识别那些”只被简单写一次、其它地方只读”的 alloca,把它们彻底消掉,直接用 SSA value 串联。
举例:_1 这个 alloca(参数 out 的栈槽)只在 entry 被参数复制写过一次,后面所有用法都是 load。mem2reg 会把它消掉,所有 mir.load v1 直接换成入口的参数 SSA value v0。
理论上 *out = 42 那三条 op 简化成:
mir.store (v0, const 42_si32) ← v0 是入口参数,直接用mem2reg 跑完后,整个函数大概会从 40 多条 op 缩到 10 几条。
这不是可选优化,而是必需的。理由:如果不跑 mem2reg,后续 lowering 会把每个 alloca 翻成 PTX 的 .local 栈访问。PTX 的 .local 是外置 DRAM,访问慢几十倍。最终 hello_constant.ptx 里这部分如果不优化,会到处都是 st.local / ld.local,kernel 性能直接崩盘。
跑完 mem2reg 后,*out = 42 在最终 PTX 里就是干干净净的三行:
ld.param.b64 %rd2, [hello_constant_param_0]; ; out 指针
cvta.to.global.u64 %rd1, %rd2; ; 转 global addrspace
st.global.b32 [%rd1], 42; ; *out = 42跟 dialect-mir 里那三条 op 一一对应——一切都对得上。
8. PHASE 6 日志在哪里
run.log 里这一步对应的日志:
[PHASE 5/9] mir-importer::run_pipeline — MIR → dialect-mir → LLVM IR → PTX functions=2 externs=0
[PHASE 6/9] translating MIR function to dialect-mir kind="kernel" name=hello_constant注意 [PHASE 6/9] 是 tracing::debug! 级别,RUST_LOG=info 看不到。想看每个函数翻译的进度:
RUST_LOG=mir_importer=debug,info cargo oxide run hello_constant[PHASE 5/9] 是 info,正常 RUST_LOG=info 就能看到。
9. 一句话总结
mir-importer 创建 pliron Context + 空 ModuleOp,逐函数翻译。每个 MIR local 在 entry 翻译成一条
mir.alloca(栈槽),每条 statement 翻译成 1-3 条 dialect-mir op。*out = 42实际生成mir.constant 42+mir.load (out_slot)+mir.store (ptr, 42)三条 op,加上 entry 那条 alloca 共四步。mem2reg 跑完之后冗余 alloca 全消掉,只剩一条 store——这不是可选优化,是 PTX 性能能跑的前提。
系列上一篇: cuda-oxide:hello-constant 拆解 05——找 #[kernel] 函数 + BFS 调用图 + stable MIR 转换