hetGPU：打破 GPU 二进制壁垒的探索

现代 GPU 生态越来越“异构”：NVIDIA、AMD、Intel 甚至 Tenstorrent 这种新架构都在同一个数据中心里出现。

但现实很残酷：GPU 二进制几乎没有跨厂商可移植性。

你给 NVIDIA 写的 CUDA（哪怕只是分发一个预编译库），在 AMD/Intel 上基本不能直接跑。原因不只是“API 不一样”，而是更底层的东西完全不兼容：

• 指令集（ISA）不同
• 执行模型不同（SIMT vs MIMD）
• 内存与一致性模型不同
• 驱动栈与工具链不同

这会带来一串非常真实的工程痛点：

• 供应商锁定：存量 CUDA 代码越大越难换硬件
• 异构集群调度很僵硬：今天 NVIDIA 忙或挂了，任务不能“拿去 AMD/Intel 顶上”
• 实时迁移几乎做不到：要容错/抢占，只能停了重跑

它的目标是：让“一个 GPU 二进制文件”可以在 NVIDIA、AMD、Intel、Tenstorrent 上跑，并且支持跨设备的运行时迁移。

如果你把它理解成“GPU 领域的 JVM / 兼容层”，就差不多对了。

1. 动机：为什么 GPU 需要“二进制兼容层”

CPU 世界其实很早就吃到“可移植字节码/虚拟化”的红利：

• Java/.NET：字节码到处跑
• 虚拟化/容器：把软件和硬件边界弱化

GPU 世界反过来是“强绑定”：

• CUDA → PTX/SASS（NVIDIA）
• HIP/ROCm → GCN/RDNA（AMD）
• oneAPI/SYCL → SPIR-V / Intel EU ISA（Intel）
• Tenstorrent → RISC-V 多核 + 向量/矩阵单元（Tensix）

你可以写 OpenCL/SYCL/HIP 做“源码可移植”，但它解决不了两个关键问题：

1. 二进制不可移植：你还是得对每个平台重新编译
2. 运行时迁移不可用：预编译库、长时间运行 kernel 的状态，都很难跨设备接续

核心矛盾其实很简单：缺乏二进制兼容性，会直接阻断 异构集群调度 和 运行中迁移/容错。

2. hetGPU 的核心想法：编译 + 运行时双管齐下

hetGPU 不是“只做一个转换器”，它是一整套系统：

• 编译器：把输入 GPU 代码编译成架构无关的 GPU IR（中间表示）
• 运行时：在加载时把 IR 动态转换成目标 GPU 的本地代码（JIT）
• 抽象层：统一线程、内存、同步语义，为不同硬件提供同一套“看起来一致”的 GPU 模型

一句话：

不直接输出 NVIDIA/AMD 的机器码，而是输出一个“可移植 IR”；真正落到哪块 GPU 上执行，由运行时决定。

2.1 一个最小例子：从 CUDA 到 hetIR（伪代码）

用一个最常见的“向量加法”来打个样。假设我们有一个 CUDA kernel：

__global__ void vadd(const float* A, const float* B, float* C, int N) {
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  if (i < N) C[i] = A[i] + B[i];
}

hetGPU 编译器不会直接生成 SASS/GCN，而是生成一个偏“架构无关”的 IR（这里用 hetIR 风格的伪代码表达）：

 tid = GET_GLOBAL_ID(0)
 pred = (tid < N)
 @pred {
   a = LD_GLOBAL_F32(A + tid * 4)
   b = LD_GLOBAL_F32(B + tid * 4)
   c = FADD_F32(a, b)
   ST_GLOBAL_F32(C + tid * 4, c)
 }
 RET

注意点：

• 里面没有写“warp=32”或“wave=64”，也没有绑定某个厂商 ISA。
• 分歧（if (i < N)）在 IR 中以谓词/结构化控制流呈现，方便后端在有/没有 SIMT 的情况下都能落地。

2.2 运行时怎么“落到具体 GPU 上”

你可以把运行时理解成一个“多后端 JIT/转换器”：

• 在 NVIDIA 上：hetIR -> PTX -> 驱动 JIT -> SASS
• 在 AMD/Intel 上：hetIR -> SPIR-V -> 驱动 JIT -> 本地 ISA
• 在 Tenstorrent 上：hetIR -> (Metalium/TT-MLIR/自定义后端) -> 目标二进制

这也解释了为什么我更愿意把它当成“一套系统工程”：只做 IR 不够，只做运行时翻译也不够。

2.3 二进制兼容性到底“兼容”什么？（别把它等同于 API 兼容）

很多人听到“GPU 二进制兼容”，第一反应是：

是不是 CUDA API 不用改，换成 AMD/Intel 也能直接跑？

但 hetGPU 讨论的重点其实更接近：交付物形态的改变。

如果你把 GPU 程序当成一个要交付的“成品”（比如一个推理算子库、一个加速组件），你会发现工程上真正难的是这三层：

• API 兼容：函数能不能调用、参数怎么传（最表层）
• 语义兼容：线程/同步/内存可见性这些“行为”是不是一致（最关键）
• 指令级兼容：NVIDIA SASS、AMD GCN/RDNA、Intel EU、Tenstorrent 指令能不能互跑（几乎别想）

hetGPU 的路子更像“GPU 版 JVM”：它尽量不让你交付任何一家厂商的设备码，而是交付一个 portable IR（hetIR），然后在目标机器上做一次“落地”。

一个很现实的发布场景：你想发 1 份库，而不是 4 份

假设你要把一个 GPU 算子发布给客户：

• 客户 A：一台 NVIDIA 服务器
• 客户 B：一台 AMD 服务器
• 客户 C：一台 Intel GPU
• 客户 D：一台 Tenstorrent

多数团队最后都会走向两条“传统路线”：

1. 让客户自己从源码编译（可移植，但交付体验差）
2. 你打包 fat binary（你编 4 份，客户按设备挑 1 份）

fat binary（胖包）的思路大概长这样：

my-kernel.pkg
  ├─ nvidia.cubin
  ├─ amd.hsaco
  ├─ intel.spv
  └─ tt.bin

它能用，但维护成本会随着设备/版本迅速爆炸。

而 hetGPU 更像是把交付物换成这一种：

my-kernel.pkg
  ├─ kernel.hetir       # “可移植的中间形态”才是主角
  └─ runtime/           # 运行时在目标机器上把 hetIR 落地到本地代码

运行时做的事情可以用一句伪代码概括：

native = JIT(hetIR, detectedGPU)
launch(native)

这就是“二进制兼容性”里最关键的那半句话：你交付的是一个统一的中间形态，而不是某家的机器码。

当然，代价也很直白：一旦把产物做成 portable，最大的难题就转移到了“语义怎么对齐”——尤其是执行模型的鸿沟（SIMT vs MIMD）。下面这节就是硬骨头。

3. 最大的坎：SIMT（warp/wave） vs MIMD（多核 RISC-V）

这也是整个问题里最“硬核”的部分之一。

3.1 NVIDIA/AMD：SIMT 是硬件自带的

NVIDIA/AMD 都是 SIMT：

• warp（NVIDIA，典型 32 threads）
• wavefront（AMD，32 或 64）

硬件会处理：

• warp 调度
• 分支分歧的 mask
• 块内同步（比如 __syncthreads()）

3.2 Tenstorrent：多核 MIMD，没有硬件 warp

Tenstorrent 的 Tensix 核心更像很多个小 CPU（RISC-V），每个核心自己跑自己的指令流。

没有 warp 调度器，意味着 CUDA 内核里那些默认假设（分歧怎么收敛、warp vote/shuffle 怎么办）都得软件补。

有两类映射思路：

1. 在一个核心内“用向量通道模拟 warp”：把 32 个线程映射到向量 lane，用 mask 处理分歧。
2. 把 warp 拆到多个核心上：核心间显式同步，靠协议模拟 warp 级协作。

这其实是在“warp 为中心”和“core 为中心”的设计之间搭一座桥。

3.3 一个“分歧怎么处理”的伪代码（向量化 warp 思路）

我们的直觉是：既然 Tenstorrent 没 warp，就在一个核心里用向量 lane 模拟 32 个线程。

下面这个伪代码抓住关键点：同一条指令流 + 掩码（mask）控制哪些 lane 生效。

// 假设一个 Tensix core 的 VPU 有 32 lanes
mask = ALL_ONES

// if (x < 0) y = -x; else y = x;
pred = (x < 0)           // pred 是每个 lane 一个 bit

// then-branch
mask_then = mask & pred
if (ANY(mask_then)) {
  y[mask_then] = -x[mask_then]
}

// else-branch
mask_else = mask & !pred
if (ANY(mask_else)) {
  y[mask_else] = x[mask_else]
}

// 汇聚点 converge
mask = mask_then | mask_else

这段代码表达的其实就是“SIMT 的分歧 + 汇聚”在 MIMD 核心上如何落地。

4. 内存模型差异：共享内存到底怎么弄

一个非常典型的差异：

• NVIDIA/AMD：块内有共享内存（shared/LDS），并且在 __syncthreads() 后有明确可见性语义
• Tenstorrent：每个核心有本地 scratchpad，跨核心共享要显式通信（DMA / mesh / global memory）

所以如果一个 CUDA kernel 强依赖 shared memory 做 32 线程协作，在 Tenstorrent 上要么：

• 把一个 block 尽量限制在一个核心里，让核心本地内存扮演 shared memory
• 或者用 global memory/显式通信去模拟 shared memory（正确但更慢）

hetGPU 的抽象层要做的就是：

给上层一个“看起来像 CUDA 的内存语义”，底层在不同设备上用不同原语去实现它。

5. 指令集差异：warp vote / shuffle / tensor 指令怎么办

这里有个经常被低估的事实：

• 很多 CUDA 内置（vote、shuffle、tensor core 指令等）在别家 ISA 里没有 1:1 对应

所以编译器/IR 必须把这些操作抽象成“统一内置”，后端再：

• 能映射就用硬件原语（NVIDIA/AMD 的 vote 类指令）
• 不能映射就用软件兜底（比如共享内存/归约通信来模拟）

这也是为什么 hetGPU 不选择“直接翻译机器码”——它更需要在一个更高层的 IR 上表达语义。

5.1 例子：没有 ballot 指令时，怎么做 vote/ballot（伪代码）

以 ballot(pred) 为例：它要返回一个 bitmask，表示一个 warp/team 里哪些线程的 pred=true。

在 NVIDIA/AMD 上可以映射到硬件 vote 指令；但在“没有现成 vote”的后端上，一个保守但可用的降级方案就是走共享缓冲区 + 汇总：

// 假设 team_size = 32
// 每个线程有 tid(0..31)

scratch[tid] = pred ? 1 : 0
barrier()

mask = 0
for i in 0..31 {
  if scratch[i] == 1 {
    mask |= (1 << i)
  }
}
barrier()
return mask

它肯定不会像硬件 vote 那么快，但它有两个好处：

• 语义清晰、跨设备一致
• 把“缺失的指令”变成“可实现的抽象原语”

5.2 例子：没有 shuffle 指令时，怎么做 warp shuffle（伪代码）

再看 shfl(val, srcLane)：想从 srcLane 那个线程拿到一个寄存器值。

硬件上通常是寄存器级交换；没有的话，也能用共享缓冲区模拟：

scratch[tid] = val
barrier()

out = scratch[srcLane]
barrier()

return out

同样的道理：性能可能差一些，但“能正确跑起来”是二进制兼容层的底线。

6. 最硬的部分：运行中状态捕获与跨 GPU 迁移

把一个 kernel 从 GPU A 停下来，换到 GPU B 接着跑，真正难点在于状态：

• 寄存器
• 程序计数器（PC）
• 分歧 mask / warp 执行状态
• shared memory
• global memory（以及一致性）

而且问题在于：

A 上的低级机器状态，很多时候没法 1:1 映射到 B。

这里还有一个非常现实的问题：底层优化会让“机器状态 → 高级状态”的映射变得多对多，几乎不可逆。

因此 hetGPU 的解决策略很像“系统工程 + 编译约束”：

• 在 IR级别插入元数据
• 明确哪些位置是“可安全捕获/恢复”的点（例如同步点）
• 用设备无关格式序列化状态

这样至少可以在“定义良好的边界”上支持迁移，而不是追求任意指令点的强行抢占。

6.1 checkpoint 到底要保存什么（一个状态结构示意）

为了跨设备恢复，你得保存“设备无关”的执行状态。大致可以抽象成下面这种结构（示意）：

Checkpoint {
  kernel_id
  grid_dim, block_dim

  // 每个线程的抽象状态（不是某个厂商 ISA 的寄存器号）
  threads: [
    { pc_ir, regs_ir[], pred_mask }
  ]

  // block 级别共享状态
  shared_mem_dump

  // 全局内存快照/增量（大头成本通常在这里）
  global_mem_regions: [ {addr, size, bytes...} ]
}

这里的关键不是“结构长什么样”，而是它必须在 IR 级别有定义，否则你只拿到某个 GPU 的机器状态，换个设备就没法解释。

6.2 为什么要强调 barrier / 安全点

想做迁移，不能指望随时把 GPU 掐断然后继续。

更靠谱的方式是：

• 在 IR 里标记可恢复点（比如 barrier、循环的某个同步点）
• 运行时在这些点触发状态捕获

原因很朴素：在 barrier 处，线程更可能对齐到一个“可解释”的状态，你保存/恢复的语义才不会崩。

6.3 例子：在 IR 里标记 safe-point（伪代码）

如果你希望“可迁移”，通常需要在 IR 层明确标记哪些位置允许暂停/恢复。一个很直观的写法是给 barrier 这类点打标签：

// ... normal instructions ...
SAFE_POINT(id=42, reason="block_barrier")
barrier()
// ... normal instructions ...

运行时触发 checkpoint 时，就可以选择“等到下一个 safe-point 再抓状态”，而不是强行在任意指令点抢占。

6.4 例子：把长 kernel 切成可恢复的 segment（伪代码）

另一种更工程化的思路是：把一个长 kernel 拆成多个段，每个段的边界都是 safe-point。

// segment 0
kernel_seg0(input, state_out) {
  // do work...
  barrier()
  dump_state(state_out)   // regs/shared/pc in IR form
}

// segment 1
kernel_seg1(state_in, output) {
  load_state(state_in)
  // continue work...
}

迁移时的控制流就变得很明确：

1. 在 GPU A 上跑 kernel_seg0，到边界时导出状态
2. 把 state_out（以及必要的 global memory）搬到 GPU B
3. 在 GPU B 上跑 kernel_seg1

它牺牲了一些“任意点抢占”的自由度，但换来的是更可解释、更可落地的恢复语义。

7. 背景对比：为什么现有方案不够

如果把现有路线粗略分一下，大概是这样：

• OpenCL/SYCL/HIP：解决源码可移植，但还是要为每个平台构建
• 胖二进制：只是把多个架构版本打包，仍然是提前准备 N 份
• API 模拟/拦截（例如 ZLUDA、Ocelot）：可行但覆盖面与正确性非常难，维护成本极高
• 检查点工具（CRIUgpu 等）：通常只针对同构/同厂商 GPU

hetGPU 想做的是把“可移植 IR + 动态转换 + 状态迁移”放在一个统一框架里。

8. 你应该怎么理解它的价值（和边界）

我觉得这个方向最大的价值不在于“它现在就能替代 CUDA”，而在于它把问题说透了：

• GPU 二进制兼容性不是一个编译器小技巧
• 它涉及 ISA、执行模型、内存语义、运行时、状态管理的协同

它能打开的想象空间是：

• 云上更灵活的 GPU 调度（不被 GPU 类型卡死）
• 故障时跨设备迁移减少重跑成本
• 新 GPU 架构的采用门槛更低（至少能先跑起来）

但你也要接受一些天然 trade-off：

• JIT/转换有成本
• 为了可迁移，可能要牺牲部分极致优化
• 对一些非常底层的特性（tensor 指令、特殊内置）可能需要“降级语义”或软件模拟

参考链接

• hetGPU 代码仓库： https://github.com/Multi-V-VM/hetGPU