论文分享-Kerncap

1 整体介绍

论文：https://arxiv.org/abs/2605.03208

Kerncap:给定“大型应用 + 目标 Kernel 名称”，自动把这个 Kernel 连同当时的参数、显存状态、二进制、源码和编译环境一起打包出来。

例如：开发者修改 Kernel 后，不再需要反复编译和运行整个 llama.cpp、LAMMPS 或 vLLM，而是直接在一个小型独立项目中执行：

修改 Kernel
    ↓
只编译该 Kernel
    ↓
恢复原始 GPU 输入
    ↓
独立执行 Kernel
    ↓
与原始输出比较

2 前置知识

一个 GPU Kernel 可能只有几十到几百行代码，但真正让它独立运行，需要同时恢复三类信息。

2.1 Kernel Definition：Kernel 定义

包括：

• Kernel 的 HSACO 二进制；
• 原始源码；
• 头文件依赖；
• 模板实例化对应的翻译单元；
• 编译宏、包含路径和优化选项；
• GPU 架构参数，例如 gfx942。

2.2 Runtime State：运行时状态

包括：

• Grid 和 Block 大小；
• 动态共享内存大小；
• Kernel 参数；
• Kernel 参数指向的显存；
• 显存中继续保存的设备指针；
• 常量内存和模块变量。

例如：

__global__ void kernel(float **weights, int *indices);

这里的 weights 并不是普通数组，而是：

Kernel 参数
    │
    ▼
设备地址 A
    │
    ▼
[A 中保存多个设备指针]
    ├── 设备地址 B → expert 0 weights
    ├── 设备地址 C → expert 1 weights
    └── 设备地址 D → expert 2 weights

只保存 weights 参数的值并不够，还必须保存 A、B、C、D 对应的全部显存。

2.3 Execution Environment：执行环境

包括：

• 如何恢复原来的虚拟地址；
• 如何重新加载 HSACO；
• 如何查找 Kernel symbol；
• 如何构造 AQL dispatch packet；
• 如何保证修改后的源码仍然使用原始编译参数；
• 如何验证新旧 Kernel 是否等价。

3 整体工作流程

3.1 HSA概念

在 AMD GPU 上，一个 HIP Kernel 从源码到真正执行，大致经过：

HIP/Triton 源码
      │
      ▼
编译器 hipcc / Triton JIT
      │
      ▼
HSACO：GPU 可以执行的代码对象
      │
      ▼
HSA Runtime：加载代码、准备参数、提交 Kernel
      │
      ▼
AQL Queue
      │
      ▼
AMD GPU 执行

所以两者的关系是：

• HSACO 是“要执行的 GPU 程序”
• HSA 层是“加载并调度这个 GPU 程序的底层运行时层”

类比 CPU 程序：

CPU 世界	AMD GPU 世界
.cpp 源码	.hip、.cu 或 Triton 源码
编译后的 .o、.so、可执行文件	HSACO
x86/ARM 机器指令	AMDGPU ISA 机器指令
操作系统加载程序	HSA Runtime 加载 HSACO

HIP 和 Triton 的上层形式差异很大：

• HIP 是提前编译的 C++ Kernel；
• Triton 是 Python 中 JIT 编译的 Kernel。

但最终二者都需要通过 HSA Runtime 向 AMD GPU 提交 AQL dispatch packet。

AQL packet 中包含：

• Kernel object；
• Grid dimensions；
• Workgroup dimensions；
• Kernarg buffer 地址；
• 动态共享内存大小；
• Completion signal。

因此，Kerncap 选择 HSA 层作为统一捕获点。

3.2 使用流程&代码

代码：https://github.com/AMDResearch/intellikit

用户层面主要使用三个命令：

kerncap profile
kerncap extract
kerncap replay

kerncap extract 内部会完成运行应用、拦截 dispatch、保存显存、查找源码和生成 reproducer 等操作。

典型提取命令类似：

kerncap extract attn_fwd \
  --cmd "python bench.py" \
  --source-dir ./src \
  --output ./isolated/attn_fwd

3.3 HIP Kernel捕获流程

Kerncap 使用 LD_PRELOAD 加载 libkerncap.so，并借助 ROCProfiler-SDK 获得 HSA Runtime 的函数指针表 HsaApiTable。然后把部分 HSA API 替换为自己的包装函数。

拦截队列

它替换 hsa_queue_create，创建 intercept queue，并为每个提交的 AQL packet 安装回调。

这样每次 Kernel dispatch 都会经过 Kerncap。

跟踪显存

它拦截：

hsa_amd_memory_pool_allocate
hsa_memory_allocate
hsa_amd_vmem_address_reserve
hsa_amd_vmem_map
free / unmap APIs

从而维护：

device virtual address → allocation size

后续显存快照就是基于这张表完成的。

跟踪 Kernel 名称

它拦截：

hsa_executable_get_symbol_by_name
hsa_executable_symbol_get_info

把底层 kernel_object 与 mangled symbol 对应起来。

捕获 HSACO

它拦截：

hsa_code_object_reader_create_from_memory
hsa_executable_load_agent_code_object

因此可以在应用加载 GPU Code Object 时，保存原始 HSACO 二进制。

发现目标 Kernel dispatch 时，不能立刻保存显存。原因是此时 Kernel 还没有执行：

dispatch callback 被调用
        │
        ├── 输入缓冲区是正确的
        └── 输出缓冲区还是旧数据

Kerncap 的做法是：

1. 创建一个新的 completion signal；
2. 用新 signal 替换 AQL packet 中原来的 signal；
3. 把 packet 提交给 GPU；
4. 等待新 signal，直到目标 Kernel 执行完；
5. 保存全部设备内存；
6. 再更新应用原来的 signal；
7. 让原应用继续执行。

因此捕获的是目标 Kernel 执行完成后的 GPU 状态。

这也意味着 Kerncap 保存的状态既可用于恢复输入，也能作为正确输出的参考。

4 核心思想：Address-Space Closure

4.1 传统方案：追踪指针图

假设 Kernel 参数是：

float **expert_weights;

传统捕获器需要：

1. 判断该参数是指针；
2. 读取它指向的显存；
3. 判断显存中的哪些 8 字节数据又是指针；
4. 继续递归读取；
5. 依赖类型信息或 DWARF 判断指针和普通整数；
6. 处理多层嵌套、数组和自定义结构体。

问题是，仅从二进制数据通常无法可靠判断：

0x7f8000000000

究竟是：

• 一个设备指针；
• 一个 64 位整数；
• 两个 32 位数；
• 某种浮点位模式。

4.2 kerncap的解决方案

不去解析哪些数据是指针，而是保存所有已跟踪的 GPU allocation，并在重放时恢复到完全相同的设备虚拟地址。

例如原应用中：

A = 0x100000
B = 0x200000
C = 0x300000

其中 A 中保存：

A[0] = 0x200000
A[1] = 0x300000

Kerncap 重放时仍然恢复为：

A → 0x100000
B → 0x200000
C → 0x300000

因此 A 中的原始字节不需要修改：

A[0] 仍指向 B
A[1] 仍指向 C

它把这一方法称为：

VA-faithful capture，保持虚拟地址忠实的捕获。

或者：

Address-space closure，地址空间闭包。

其本质是：

保存指针图

转化为：

保存整个已知地址空间，并恢复原始地址

这是一种典型的“用空间换通用性”方法。它可能保存大量与目标 Kernel 无关的显存，但避免了复杂且不可靠的指针分析。

4.3 vLLM的例子

vLLM 的 fused_moe_kernel 会通过多层指针访问不同专家的权重。

论文中 Kerncap 捕获了：

• 185 个显存区域；
• 约 30 GB 状态；
• 约 24 GB 专家权重；
• 其余为 activation 和 workspace。

这些权重并不是全部直接作为 Kernel 参数传递，而是通过 MoE 路由结构中的指针间接访问。

4.4 缺陷

方法不是完全通用的，论文明确列出几种失败情况。

Host Pointer

如果 Kernel 参数中存在：

• Host-mapped memory；
• Unified memory 中的主机地址；
• mmap 文件地址；
• CPU 地址；

那么仅恢复设备 VMEM 不够。

未被 HSA API 跟踪的地址

如果某个设备指针对应的分配没有经过 Kerncap 拦截的 HSA API，它不在 allocation map 中，也就不会被保存。

跨进程的 Host ASLR

设备虚拟地址可以通过 HSA VMEM API请求原地址，但主机虚拟地址受 ASLR 影响，通常无法保证在新进程中相同。

4.5 常量内存和模块变量

只拦截显存分配 API 还不够。

例如 Kokkos 会使用：

kokkos_impl_hip_constant_memory_buffer

这类数据可能通过 hipMemcpyToSymbol 写入 __constant__ 或模块变量。

这些内存并不一定出现在普通 allocation map 中。

Kerncap 因此还会：

1. 枚举 HSACO 中的 HSA_SYMBOL_KIND_VARIABLE；
2. 保存变量内容到 module_vars/；
3. 重放时，在 Kernel dispatch 前恢复；
4. 使用 HSACO SHA-256 区分同名模块变量。

5 注意事项

5.1 固定配置

Triton 的 @triton.autotune 会在多个配置中选择最快者，例如：

BLOCK_M = 64
BLOCK_N = 64
num_warps = 4
num_stages = 2

不同 tile size 不只影响性能，也可能改变浮点运算顺序。

例如 reduction 被划分成不同 tile 后：

(a + b) + c

可能变成：

a + (b + c)

在 FP16/FP32 中两者不一定完全相同。

论文对 Flash Attention 的 attn_fwd 进行了实验：

• 最快配置和第二快配置只相差约 7.7% 性能；
• 但约 11.3% 的输出元素发生变化；

因此，重新运行提取后的 Triton Kernel 时，如果再次让 autotuner 搜索，可能选择不同配置，从而导致输出与捕获时不同。

Kerncap 的方案是生成 tuning-pinned reproducer：

kernel.fn[grid](..., BLOCK_M=..., BLOCK_N=..., num_warps=...)

绕过 autotuner wrapper，直接绑定捕获时的 winning configuration。

论文将其描述为保持 Triton JIT Kernel 的“数值契约”。

总结：对一个核函数内部的配置，kerncap会选择一个最优的配置保存下来

5.2 自动寻找HIP源码

捕获 HSACO 和显存只能实现二进制重放；如果想修改 Kernel，还需要找到源代码。

HIP 源码发现采用多阶段策略。

1. 优先使用 DWARF

当项目存在：

compile_commands.json
+
带调试信息的 object file

Kerncap 会：

1. 从 dispatch.json 获得精确 mangled symbol；
2. 使用 nm 查找哪个 .o 文件包含该 symbol；
3. 使用 llvm-dwarfdump 读取 DWARF line table；
4. 找到参与该翻译单元的源码；
5. 根据 --source-dir 过滤文件；
6. 从 compile_commands.json 恢复编译命令。

这一方案适合 Kokkos 等框架生成的复杂 Kernel，因为用户源码中可能根本找不到直接的 __global__ 定义。

2. Grep 作为后备

如果没有 DWARF 或编译数据库，则执行两阶段搜索：

第一阶段严格查找：

__global__ ... kernel_name(...)

第二阶段仅搜索函数基础名称。

找到文件后，再递归跟踪本地 #include，论文实现中最多跟踪五层。

3. 处理模板实例化

像 llama.cpp 中可能存在：

mmq-instance-1.cu
mmq-instance-2.cu
mmq-instance-3.cu

它们都包含相同模板头文件，但只实例化了不同类型或参数组合。

仅通过字符串搜索无法判断目标 Kernel 来自哪个翻译单元。

Kerncap 使用精确 mangled symbol，通过 nm 检查 object file，从多个候选中选出真正产生该 Kernel 的翻译单元。

6 实验与结果

6.1 benchmark

论文在三种 AMD GPU 架构上进行了测试：

GPU	架构
AMD MI300X	CDNA3，gfx942
AMD MI210	CDNA2，gfx90a
Radeon PRO W7900	RDNA3，gfx1100

软件环境为 ROCm 7.2.0、RHEL 9.6。

覆盖六个工作负载：

工作负载	类型	目标 Kernel
llama.cpp	HIP	mul_mat_vec_q
LAMMPS/Kokkos	HIP	TagPairEAMKernelC
rocBLAS/Tensile	HIP HSACO	GEMM Kernel
Flash Attention 2	Triton	attn_fwd
vLLM MoE	Triton	fused_moe_kernel
torch.compile	Triton/Inductor	fused ReLU Kernel

这些 workload 覆盖：

• 手写 Kernel；
• 模板生成 Kernel；
• Kokkos；
• Tensile；
• Triton；
• PyTorch Inductor；
• 传统 HPC；
• LLM 推理；
• 152 MB 到约 30 GB 的显存状态。

6.2 结果

Kerncap 在论文测试的所有 workload/架构组合上都完成了提取和验证。

一些代表性结果：

工作负载	GPU	区域数	快照大小	总提取时间
llama.cpp	gfx942	21	12,233 MB	15.6 s
LAMMPS	gfx942	84	8,449 MB	25.0 s
rocBLAS	gfx1100	10	152 MB	4.4 s
Flash Attention	gfx942	13	178 MB	9.1 s
vLLM	gfx942	185	30,074 MB	60.9 s
torch.compile	gfx942	12	227 MB	9.4 s

vLLM 在 gfx90a 上约需 85.2 秒，是表中时间最长的情况，但它保存了接近 30 GB 的状态。

这里要注意：Kerncap 不追求最小快照，而追求“足够完整、无需理解指针”的快照。

6.3 捕获开销

论文将开销分成两部分。

6.3.1 Interception Tax

即每个 HSA 调用和 dispatch 回调经过 Kerncap 所带来的持续开销。

对于较大的 HIP 应用：

• CDNA 上通常不超过约 1.2 倍；
• RDNA 上较大 workload 通常不超过约 1.4 倍。

rocBLAS 微基准的比例较高，是因为基线只有几秒钟，固定回调开销被放大。

6.3.2 Snapshot Cost

目标 Kernel 捕获时的一次性显存复制成本。

它主要取决于：

显存快照大小 ÷ 实际 D2H 与磁盘吞吐

而不是取决于应用中 Kernel 的总数量。

论文在 gfx942 上观察到约 1.7 GB/s 的综合快照吞吐。

因此：

• 8 GB 快照大约几秒；
• 12 GB 快照约 7 秒；
• 30 GB 快照约十几秒到几十秒。

对于运行几分钟的大型应用，一次性捕获成本所占比例并不高。

7 优化案例

llama.cpp 优化案例

目标 Kernel：

mul_mat_vec_q

它约占 llama.cpp GPU 总时间的 8%，调用约 7,740 次。

原始性能

独立重放十次后，平均时间：

$34.5 μs$

优化方法

进行了两项代码变换：

1. 把 if (use_gate) 从高度展开的内层循环中移到循环外；
2. 使用模板和 if constexpr 强制编译期分支，并内联 dot-product 函数。

效果包括：

• 消除内层运行时分支；
• 避免寄存器 spilling；
• 生成完整展开的无分支循环。

优化后性能

$28.3 μs$

即：

$\frac{34.5}{28.3}\approx1.22\times$

执行时间下降约 18%。

修改后的 Kernel 通过了全部 21 个显存区域的字节级一致性验证。

8 迭代速度提升

llama.cpp

传统流程：

完整增量构建：128 s
应用运行：8.3 s
合计：约 136 s

Kerncap：

单 Kernel 编译：18.3 s
独立 replay：6.9 s
合计：约 25 s

快速内循环提升：

$136/25≈5.4×$

LAMMPS

传统内循环约 74 秒，Kerncap 约 18 秒：

$约4.1×$

包含严格验证

llama.cpp 严格字节验证需要约 129.4 秒，而传统应用级 ROCm Compute Profiler 流程约需 2,072 秒。

完整的：

编辑
编译
重放
严格验证

Kerncap 约 162 秒，传统流程约 2,208 秒：

$约13.6×$

论文标题和摘要中的 13.6 倍，指的是这一包含完整验证的端到端工作流，而不是单纯 Kernel 执行速度。

9 论文创新点

创新一：统一 HIP 和 Triton 的 HSA 捕获

HIP 和 Triton 在上层差异很大，但底层都经过 HSA。

Kerncap采用：

统一 HSA 捕获
+
Triton Python 元数据补充
+
语言特定的 Reproducer

既共享底层机制，又避免强行把 HIP 和 Triton处理成完全相同的语言。

创新二：VA-Faithful Address-Space Closure

这是最核心的系统创新。

它不追踪指针，而是保证所有设备虚拟地址不变，从而自动保存任意深度的设备指针关系。

创新三：运行时信息辅助源码发现

普通静态工具只知道源码和编译数据库，不知道实际执行的是哪个模板实例。

Kerncap 将：

运行时 mangled symbol
+
compile_commands.json
+
DWARF
+
nm

组合起来，解决模板翻译单元歧义。

创新四：Clang VFS 隔离编译

不修改原项目、不重写复杂编译命令，仅通过 VFS Overlay 替换目标源码。

创新五：固定 Triton Autotuner 的数值契约

论文不仅考虑“能运行”，还注意到 autotune 配置变化会导致浮点输出变化，因此保存并固定 winning configuration。

10 论文局限性

1.一次只捕获一个 Kernel Dispatch

Kerncap不能直接捕获：

Kernel A → Kernel B → Kernel C

作为一个联合执行单元。

虽然捕获 B 时，A 已经执行完，其输出会被包含在快照中，但无法研究 A 和 B 的联合变换、跨 Kernel fusion 或流水线调度。

2.只支持单 GPU

多 GPU 应用需要：

• 区分多个 HSA Agent；
• 同步多个设备的捕获时间；
• 保存 GPU 间通信状态；
• 恢复 RCCL 或 P2P 关系。

当前版本没有解决。

3.不捕获 Host State

不保存：

• CPU 内存；
• 文件描述符；
• 主机映射缓冲区；
• 环境变量；
• 网络通信状态。

因此依赖 host-mapped memory 的 Kernel 可能无法重放。

4.快照可能非常大

Address-space closure 很稳健，但代价是：

它捕获所有被跟踪的 allocation，而不是只捕获目标 Kernel 真正读取的数据。

vLLM 单个 Kernel 的 reproducer 接近 30 GB，这对大规模 Agent 调优会产生明显的：

• 磁盘占用；
• 加载时间；
• 初始化成本；
• 多候选并发运行压力。

5.源码发现不是完全可靠

没有 compile_commands.json 和 DWARF 时，Kerncap依赖 grep 和 include tracing。

对于：

• 宏生成 Kernel；
• 复杂 CMake 生成文件；
• 计算式 include；
• 动态生成源码；
• 多层代码生成框架；

可能找不到可编辑源码。

6.只验证捕获的运行上下文

Kerncap证明的是：

修改后的 Kernel 在这一组捕获输入上，与基线一致。

它并不能自动证明：

修改后的 Kernel 对所有输入、所有 shape、所有边界情况均正确。

因此，候选 Kernel 可能过拟合单次 dispatch。这一点对自动调优和 LLM Agent 尤其重要。

7.Autotuner 配置不具备跨架构可移植性

MI300X 上的最佳：

BLOCK_M / BLOCK_N / num_warps

不一定适用于 MI210 或 RDNA。

Kerncap固定配置是为了忠实重放，而不是实现跨 GPU 调优。

论文提出的未来方向包括：

• 多 Kernel 联合捕获；
• 多 GPU 支持；
• 与自动调优系统闭环集成；
• CUDA/CUPTI 支持；
• 建立持久化 Kernel 数据库，用于 ROCm 版本回归测试