降级流水线#

前面章节从两端构建了 IR：MIR 导入器将 Stable MIR 翻译为 dialect-mir，而 Pliron 方言描述了 dialect-mir、dialect-llvm 和 dialect-nvvm。本章介绍它们之间的桥梁——将 Rust 风格的 IR 转换为 LLVM 能够实际编译的东西的 pass。

如果你了解 Rust 类型，你即将发现 LLVM 从未听说过其中的多少种。

降级的含义#

dialect-mir 说着 Rust 的语言。它了解元组、枚举、切片、检查算术和 GPU 地址空间。LLVM IR 对这些一无所知。它具有扁平的整型和浮点类型、getelementptr、PHI 节点，以及对任何多于一个抽象层次的普遍怀疑。

降级是将每个 dialect-mir 操作逐一替换为等价的 dialect-llvm 操作序列的过程，直到不再有 dialect-mir 操作剩余。元组变成匿名结构体。切片变成指针-长度对。检查加法变成 LLVM 溢出内建函数后跟一个 extract。每个 Rust 概念都被展平为 LLVM 可以消化的东西。

完成这一切的 pass 位于 crates/mir-lower/，并使用 pliron 的 DialectConversion 框架。它是流水线中最大的单个转换，本章剩余部分将介绍它是如何工作的。

DialectConversion -- 降级框架#

降级使用 pliron 的 DialectConversion + DialectConversionRewriter，而非手动的遍历替换 pass。该框架自动处理 IR 遍历、def-before-use 排序、类型转换和块参数修补。

工作原理#

每个 dialect-mir 和 dialect-nvvm 操作通过 MirToLlvmConversion 操作接口声明如何降级自身（定义于 conversion_interface.rs）。该接口只有一个方法：convert(ctx, rewriter, op, operands_info)。每个操作的实现位于 convert/interface_impls.rs，它分发到按类别组织的转换函数。

对于模块中的每个 MirFuncOp，convert_func（在 lowering.rs 中）：

  1. 创建 LLVM 函数，其签名取决于该函数是否为内核入口点。非内核函数将每个聚合参数展平为其标量字段，以符合内部 CUDA ABI。内核入口点保持切片展平（(ptr, len)），但将结构体和闭包作为单个 byval 值传递，因为主机启动器将整个聚合作为一个数据包槽位推入（参见下方的参数标量化）。

  2. 传播 GPU 元数据（gpu_kernel、maxntid、cluster_dim_*）。

  3. 使用 inline_region 将 MIR 块移入 LLVM 函数。原始块被保留——无需手动块映射。

  4. 构建入口前导代码，将每个参数的 MIR 级别值从到达 LLVM 级别的任何形式重新构造出来（对于内部调用和切片：展平标量的 insertvalue；对于内核边界的 byval 聚合：直接传递）。

  5. 运行 DialectConversion，遍历每个 MIR 操作并调用其 MirToLlvmConversion::rewrite 实现将其替换为 LLVM 操作。

转换模块#

转换函数按类别组织为模块：

类别

模块

处理内容

算术

convert/ops/arithmetic.rs

add→add、sub→sub、checked_add→add+extractvalue

内存

convert/ops/memory.rs

mir.load→load、mir.store→store、shared_alloc→global + addrspacecast

控制流

convert/ops/control_flow.rs

mir.goto→br、mir.cond_br→cond_br、mir.return→return

聚合

convert/ops/aggregate.rs

结构体/元组字段访问 → GEP 或 extractvalue/insertvalue

类型转换

convert/ops/cast.rs

IntToInt→zext/sext/trunc、FloatToFloat→fpext/fptrunc 等

调用

convert/ops/call.rs

mir.call→call，带参数展平和 :: 到 __ 的名称转换

GPU 内建函数

convert/intrinsics/*.rs

NVVM 操作 → LLVM 内建函数调用或内联 PTX

常量

convert/ops/constants.rs

mir.constant→llvm.constant

框架自动分发到这些模块——每个操作的 MirToLlvmConversion impl 调用正确的转换函数。复杂性存在于每个转换器内部，即 Rust 语义与 LLVM 实际交汇的地方。

类型转换#

在操作可以被转换之前，它们的类型必须被转换。LLVM 的类型系统有意比 Rust 的更简单——整数上没有符号性，没有元组，没有枚举，没有胖指针。一切都必须被展平。

MIR 类型

LLVM 类型

备注

IntegerType(32, Unsigned)

IntegerType(32, Signless)

LLVM 整数不携带符号信息——符号性在操作上，而非类型上

MirTupleType<i32, f32>

{ i32, float }

元组变成匿名结构体

MirSliceType<f32>

{ ptr, i64 }

胖指针分解——指针 + 长度

MirStructType

{ fields..., [N x i8] }

显式的填充数组以匹配 rustc 的布局

MirPtrType<f32, addrspace:3>

ptr addrspace(3)

不透明指针，保留地址空间

MirArrayType<f32, 256>

[256 x float]

直接映射——数组对 LLVM 来说足够简单

MirEnumType

{ discriminant, [M x i8] }

判别值 + 按最大变体尺寸调整的有效载荷

整数符号性的情况值得强调。在 Rust 中，i32 和 u32 是不同的类型。在 LLVM 中，两者都只是 i32。符号信息转移到操作上：有符号小于比较是 icmp slt，无符号的是 icmp ult。类型转换器丢弃符号性，操作转换器在发出比较和除法指令时重新获取它。

参数标量化#

内核入口点需要特殊处理。CUDA 驱动不理解 Rust 胖指针，因此 &[f32] 必须在 ABI 两侧分别以指针和长度到达。另一方面，按值传递的结构体和闭包确实匹配单个主机数据包槽位，因此内核入口将其作为单个 byval .param 接收——否则设备会期望 N 个展平参数，而主机只会推送一个，导致后面的每个切片参数错位。

因此降级 pass 区分内核入口规则和内部调用规则：

MIR 内核参数

LLVM 内核签名

&[f32]

ptr addrspace(1) %ptr, i64 %len

T（标量）

直接传递

struct { a, b }

单个 byval {a, b} 值

闭包（N 个捕获）

单个 byval 闭包结构体值

零大小聚合

丢弃（无 LLVM 参数，无主机数据包槽位）

切片的情况仍然在入口块中使用经典的从展平重新构造模式：

MIR:  fn kernel(slice: &[f32])
      → 入口参数: %slice : MirSliceType

LLVM: fn kernel(ptr addrspace(1) %ptr, i64 %len)
      → 入口块重新构造:
          %slice  = insertvalue {ptr, i64} undef, %ptr, 0
          %slice2 = insertvalue {ptr, i64} %slice, %len, 1

结构体/闭包的情况跳过重新构造——byval 值已经是正确的形状——函数的其余部分就像没有什么特殊的事情发生过一样看到它。内部 device-to-device 调用保持一如既往的展平聚合，因此这条规则的成本仅存在于内核边界。

有趣的转换#

大多数转换是直接的：整数上的 mir.add 变为 llvm.add，mir.load 变为 llvm.load，mir.goto 变为 llvm.br。有趣的案例是那些单个 MIR 操作展开为多个 LLVM 操作，或者 GPU 特定的关注点完全改变转换的情况。

检查算术#

在 debug 构建中，Rust 检查每个整数算术操作的溢出。MIR 使用 mir.checked_add 等返回 (result, overflow_flag) 元组的操作来建模。LLVM 没有这样的概念，但它有溢出内建函数：

MIR:  %result = mir.checked_add %a, %b : i32  → mir.tuple<i32, bool>

LLVM: %sum = add i32 %a, %b
      %overflow = extractvalue {i32, i1} @llvm.sadd.with.overflow.i32(%a, %b), 1

溢出标志馈入一个 assert，该 assert 已由 MIR 导入器降级为指向 unreachable 块的条件分支。在 GPU 上，这实际上意味着：如果你设法触发整数溢出，内核会陷入陷阱。这不是最优雅的错误处理方式，但 CUDA 工具链目前不支持栈展开。

共享内存#

CUDA 中的共享内存是块作用域的 SRAM——快速、小巧、静态声明。在 dialect-mir 中，它是一个 mir.shared_alloc 操作。在 LLVM IR 中，共享内存必须是地址空间 3 中的模块级全局变量：

MIR:  %shmem = mir.shared_alloc : mir.array<f32, 256>

LLVM: @shmem_0 = addrspace(3) global [256 x float] zeroinitializer
      %ptr = addrspacecast [256 x float] addrspace(3)* @shmem_0 to ptr

addrspacecast 产生一个通用指针，函数的其余部分可以使用它而无需关心地址空间。LLVM 中的 NVPTX 后端处理其余部分——它知道 addrspace(3) 意味着共享内存并生成适当的 st.shared / ld.shared 指令。

枚举降级#

Rust 枚举在代数上很丰富。LLVM 没有带标签联合的概念。降级 pass 通过将枚举表示为具有两个字段的结构体来弥合这一差距：一个判别值（告诉你哪个变体是活动的）和一个按最大变体大小调整的有效载荷区域：

MIR:  %opt = mir.construct_enum "Some", (%val) : mir.enum<"Option_i32">

LLVM: %tmp = insertvalue { i8, [4 x i8] } zeroinitializer, i8 1, 0
      %result = insertvalue { i8, [4 x i8] } %tmp, <val into payload area>

判别值是 i8 1，因为 Some 是 Option 的变体 1。有效载荷是 [4 x i8]——四个字节，足够容纳一个 i32。变体访问反向工作：读取判别值，在其上分支，然后 extractvalue 有效载荷并 bitcast 到预期类型。

这不够优雅，但这正是 C 编译器数十年处理带标签联合的方式。LLVM 的优化器非常擅长清理冗余的 insertvalue/extractvalue 链。

GPU 内建函数转换#

dialect-nvvm 操作——线程索引、warp shuffle、barrier、TMA 批量拷贝——不会被降级为通用的 dialect-llvm 操作。它们被降级为 LLVM 内建函数调用或内联 PTX 汇编，具体取决于 LLVM 是否有该操作的内置内建函数。

策略 1：LLVM 内建函数调用#

对于 LLVM 已经提供目标特定内建函数的操作，转换发出对该内建函数的 call：

nvvm.read_ptx_sreg_tid_x
  → call i32 @llvm_nvvm_read_ptx_sreg_tid_x()

nvvm.shfl_sync_bfly_i32
  → call i32 @llvm_nvvm_shfl_sync_bfly_i32(i32 -1, i32 %val, i32 %mask, i32 31)

注意 warp shuffle：面向用户的 cuda_device API 接受两个参数（值和 lane mask），但 LLVM 内建函数接受四个（membermask、value、delta、clamp）。降级 pass 填充缺失的参数——membermask = -1（所有 lane）和 clamp = 31（完整的 warp 宽度）——因此用户无需考虑它们。

策略 2：内联 PTX 汇编#

较新的 GPU 指令通常缺少 LLVM 内建函数。对于这些，降级 pass 使用 LLVM 的 asm 语法发出内联 PTX 汇编：

nvvm.wgmma_fence_sync
  → call void asm sideeffect convergent "wgmma.fence.sync.aligned;", ""()

nvvm.mbarrier_arrive
  → call i64 asm sideeffect convergent "mbarrier.arrive.shared.b64 $0, [$1];", "=l,r"(ptr %bar)

这里的 convergent 属性至关重要。它告诉 LLVM："不要跨控制流移动、复制或推测此指令。"如果没有它，LLVM 可能会将 barrier 提升出条件分支，或将 warp-level 指令下沉过同步点，导致 GPU 挂起或计算出垃圾——这两种结果都不会产生有用的错误消息。

块参数到 PHI 节点#

Pliron IR（类 MLIR）使用块参数进行基本块之间的值流动。LLVM 使用 PHI 节点。它们表达相同的概念——"此值来自不同的前驱"——但语法不同，导出步骤需要真正的转换，而不仅仅是美化打印。

Pliron 风格（块参数）：

^loop_header(%sum: f32, %i: i64):
    ...
    br ^loop_header(%new_sum, %new_i)

LLVM IR 风格（PHI 节点）：

loop_header:
    %sum = phi float [ 0.0, %preheader ], [ %new_sum, %body ]
    %i = phi i64 [ 0, %preheader ], [ %new_i, %body ]

导出器通过两遍方法处理此转换：

  1. 预 pass：命名每个值。 在发出任何代码之前，导出器遍历所有块并为每个值分配顺序的 SSA 名称（%v0、%v1、...）。这很关键，因为 PHI 节点可以引用在列表中较后出现的块中的值——循环回边在文本中向前指向，但在控制流中向后指向。如果不预命名，这些引用将是未定义的。

  2. 构建前驱映射。 对于每个块，导出器通过检查函数中的每条分支指令来收集 (predecessor_block, values_passed) 对。

  3. 发出 PHI 节点。 在每个非入口块的入口处，导出器为每个块参数发出一个 PHI 节点，填充来自前驱映射的值和前驱标签。

预 pass 是微妙的环节。考虑一个循环：循环头部的 PHI 引用了来自循环体的 %new_sum，但循环体在文本输出中出现在头部之后。如果我们在发出过程中即时分配名称，%new_sum 还没有名称。预 pass 通过预先命名所有内容消除了这一问题。

符号名称清理#

函数名称经过多个阶段，每个阶段施加自己的约束：

rustc_public (FQDN)            helper_fn::cuda_oxide_device_<hash>_vecadd
  ↓ body.rs (:: → __)
dialect-mir                    helper_fn__cuda_oxide_device_<hash>_vecadd
  ↓ call.rs (:: → __)
dialect-llvm                   helper_fn__cuda_oxide_device_<hash>_vecadd
  ↓ export.rs (去除前缀)
文本 LLVM IR                   @vecadd
  ↓ llc
PTX                            vecadd

在此路径上发生三次转换：

  1. :: 到 __——body.rs（函数定义）和 call.rs（调用目标）都将 Rust 路径分隔符替换为双下划线以产生有效的 pliron/LLVM 标识符。由于双方应用相同的转换，定义和调用点匹配。

  2. 设备前缀剥离——export.rs 通过 reserved_oxide_symbols::device_base_name 从 #[device] 函数名称中剥离保留的 cuda_oxide_device_<hash>_ 前缀（以及任何前面的 FQDN crate 前缀）。此前缀用于 MIR 级别检测，但不应出现在最终的 LLVM IR、PTX 或 LTOIR 输出中。

  3. 设备 extern 前缀剥离——对于 #[device] unsafe extern "C" 函数，call.rs 通过 reserved_oxide_symbols::device_extern_base_name 剥离 cuda_oxide_device_extern_<hash>_ 前缀，以便 LLVM IR 引用外部 LTOIR（例如 CCCL 库）导出的原始符号名称。

备注

当框架升级时，此手动清理将被 pliron 的 Legaliser 取代。Legaliser 系统地处理 :: 到 _ 的转换和冲突检测。

PTX 生成#

在 dialect-llvm 被导出为文本 .ll 文件后，最后一步是调用 llc——LLVM 的静态编译器——生成 PTX 汇编：

llc -march=nvptx64 -mcpu=sm_90 kernel.ll -o kernel.ptx

目标选择#

流水线按以下顺序探测 PATH 上的 llc。LLVM 21 是最低要求——较早的版本拒绝 cuda-oxide 发出的 TMA / tcgen05 / WGMMA 内建函数签名。

优先级

llc 版本

目标

PTX 版本

第 1

llc-22

sm_100a（Blackwell DC）

PTX 8.x

第 2

llc-21

sm_100 / sm_120

PTX 8.x

如果两者都不可用，流水线会以清晰的错误消息失败。你可以通过设置 CUDA_OXIDE_LLC=/path/to/llc 选择特定的（可能较旧的）二进制文件，但只有简单的内核保证能在 LLVM 20 及以下版本上编译。

如果选定的目标与物理 GPU 不匹配，CUDA 驱动会在加载时 JIT 编译 PTX。首次启动约花费 30ms，驱动进行翻译；后续启动使用缓存的二进制文件。实践中你很少注意到——JIT 很快，缓存在多次运行之间持久。

CUDA_OXIDE_TARGET 环境变量覆盖自动检测，适用于你需要特定目标的情况。例如，sm_100a 启用了在通用 sm_100 目标下不可用的 Blackwell 特定的 tcgen05 功能。

cargo oxide run 在后端基于功能的默认值之上添加了第二层自动检测：当 --arch 和 CUDA_OXIDE_TARGET 都未设置时，它会查询 CUDA 设备 0 的计算能力并将其转发给后端，使得生成的模块保证能在本地 GPU 上加载。完整的优先级是 --arch > CUDA_OXIDE_TARGET > 主机 CC（仅 run） > 后端基于功能的默认值。cargo oxide build 和 cargo oxide pipeline 有意跳过主机 CC 步骤，以便它们仍可用于交叉编译。

备注

为什么需要 LLVM 21？tma_copy、gemm_sol 和 tcgen05_matmul 使用的 2-D 批量 TMA 加载内建函数在 LLVM 21 中获得了带有 addrspace(7) 和 cta_group 参数的 10 操作数形式。较旧的 llc 版本会以 Intrinsic has incorrect argument type! 拒绝它。我们选择 21 作为最低要求，而非为每个 LLVM 版本维护单独的内建函数发射器。

总结#

以下是 lower_mir_to_llvm 处理模块时的完整事件序列：

1. 注册 `dialect-llvm` 类型和操作
2. 对于模块中的每个 MirFuncOp：
   a. 创建带有扁平类型签名的 `llvm.func`
   b. inline_region：将 `dialect-mir` 块移入 LLVM 函数
   c. 构建入口前导代码（从扁平参数重新构造聚合体）
   d. 运行 DialectConversion：
       ├── 按 def-before-use 顺序遍历每个 `dialect-mir`/`dialect-nvvm` 操作
       ├── 为每个操作调用 MirToLlvmConversion::rewrite
       ├── 转换器通过 DialectConversionRewriter 发出 `dialect-llvm` 操作
       └── 框架自动修补块参数类型
3. 将 `dialect-llvm` 导出为文本 LLVM IR (.ll)（含 PHI 节点转换）
4. 调用 llc 生成 .ptx

在第 4 步之后，你有了一个 CUDA 驱动可以加载和执行的 .ptx 文件。从 mir.checked_add 到 add.s32 的旅程完成了。

降级流水线将 Rust 风格的 IR 转换为 GPU 就绪的 LLVM IR。关于添加新 GPU 操作的实际操作指导，请参见 添加新的内建函数。