从 cubin 到 kern：引擎对模型的感知面可以压缩到什么程度

Sep 4, 2026

从 cubin 的定义出发，再到如何 dump 一个程序的 cubin，最后落到 kern 的 manifest 设计：引擎对模型的感知面可以压缩到什么程度。

1. TL;DR

cubin 是单架构的 ELF 设备二进制，其 entry symbol 经 demangle 后能直接读出 kernel 的形参列表，这是”只拿二进制就能执行”的根基。
借助 CUDA 的 profiler 注入机制，可以在不改目标程序的前提下抓到它实际加载的 module 和实际发生的 launch。一次 vLLM 服务的 bring-up 中实际用到的 cubin 约 90 个、launch 约 7 万次。
Driver API 执行一个 kernel 只需要 module、symbol、grid/block 和一段参数内存。参数布局由 kernel image 自身元数据定义，launch 端不需要感知实现。
kern 把引擎需要感知的东西收敛为四层声明：state 布局、buffer、kernel 来源与签名、调用序列。

2. cubin 是什么

2.1 三种产物

根据 CUDA Binary Utilities 与 NVCC 规范，不同阶段的产物有明确的职责划分：

后缀 / 格式	官方定位	目标架构	表现形式
`.ptx`	虚拟指令集（ISA）的中间汇编	Virtual (`compute_*`)	纯文本
`.cubin`	针对单一 GPU 架构的设备二进制	Real (`sm_*`)	ELF 格式二进制
`.fatbin`	多版本设备码集合容器	混合（支持多虚拟/真实架构）	包装容器

2.2 编译轨迹与打包机制

双架构编译模型：nvcc 设备端编译依赖两级架构定义。前端编译器先将 C/C++ 源码转译为虚拟架构的 PTX 汇编（如 compute_90），再经由 ptxas 将 PTX 汇编为目标硬件 ISA 的 cubin。
参数控制路径：
- 指定 -cubin：仅生成单架构设备端 .cubin 文件，主动丢弃 host 端代码。
- 默认编译流：将编译出的 cubin 及/或 PTX 打包进 fatbin，最终嵌入 host 端二进制，由运行时根据当前执行环境自动选择最匹配的镜像。

3. 一个 VecAdd 的三个 level

我们以 Programming Guide 里的 VecAdd（C[i] = A[i] + B[i]，设备端无 printf）为例，加一个 main 让默认 nvcc 路径能吐出 host 二进制：

// Programming Guide VecAdd kernel, with a small host so the default
// nvcc path produces a host binary.
#include <stdio.h>

__global__ void VecAdd(float *A, float *B, float *C) {
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main() {
    const int N = 4;
    float h_A[4] = {1.f, 2.f, 3.f, 4.f};
    float h_B[4] = {10.f, 20.f, 30.f, 40.f};
    float h_C[4];

    float *A, *B, *C;
    cudaMalloc(&A, N * sizeof(float));
    cudaMalloc(&B, N * sizeof(float));
    cudaMalloc(&C, N * sizeof(float));
    cudaMemcpy(A, h_A, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(B, h_B, N * sizeof(float), cudaMemcpyHostToDevice);

    VecAdd<<<1, N>>>(A, B, C);
    cudaDeviceSynchronize();
    cudaMemcpy(h_C, C, N * sizeof(float), cudaMemcpyDeviceToHost);

    int ok = 1;
    for (int i = 0; i < N; i++) {
        if (h_C[i] != h_A[i] + h_B[i]) {
            ok = 0;
        }
    }
    cudaFree(A);
    cudaFree(B);
    cudaFree(C);
    return ok ? 0 : 1;
}

本机环境：CUDA 13.3（V13.3.33），-arch=native，RTX 5070 Ti → sm_120。三条 nvcc 路径：

nvcc -cubin -arch=native -o kernel.cubin kernel.cu   # 丢掉 host 代码，只写设备端 .cubin
nvcc -ptx   -arch=native -o kernel.ptx   kernel.cu   # 虚拟 ISA 文本
nvcc        -arch=native -o kernel       kernel.cu   # 默认：cubin/PTX → fatbin → 嵌入 host 二进制

3.1 PTX level

先看第一个 level，也就是 PTX。文件头是这样的：

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-37862127
// Cuda compilation tools, release 13.3, V13.3.33
// Based on NVVM 23.0.0
//

.version 9.3
.target sm_120
.address_size 64

分别是 PTX 语言版本、target 架构、address size（指针按 64 位算）。

然后就是 kernel 本体的 PTX。这些指令都能在 PTX ISA 文档里找到对应的声明，这里就不一一展开：

    // .globl    _Z6VecAddPfS_S_

.visible .entry _Z6VecAddPfS_S_(
    .param .u64 .ptr .align 1 _Z6VecAddPfS_S__param_0,
    .param .u64 .ptr .align 1 _Z6VecAddPfS_S__param_1,
    .param .u64 .ptr .align 1 _Z6VecAddPfS_S__param_2
)
{
    .reg .b32     %r<5>;
    .reg .b64     %rd<11>;

    ld.param.b64     %rd1, [_Z6VecAddPfS_S__param_0];
    ld.param.b64     %rd2, [_Z6VecAddPfS_S__param_1];
    ld.param.b64     %rd3, [_Z6VecAddPfS_S__param_2];
    cvta.to.global.u64     %rd4, %rd3;
    cvta.to.global.u64     %rd5, %rd2;
    cvta.to.global.u64     %rd6, %rd1;
    mov.u32     %r1, %tid.x;
    mul.wide.u32     %rd7, %r1, 4;
    add.s64     %rd8, %rd6, %rd7;
    ld.global.b32     %r2, [%rd8];
    add.s64     %rd9, %rd5, %rd7;
    ld.global.b32     %r3, [%rd9];
    add.f32     %r4, %r2, %r3;
    add.s64     %rd10, %rd4, %rd7;
    st.global.b32     [%rd10], %r4;
    ret;
}

3.2 cubin level

下一个 level 就到了 kernel.cubin。这玩意儿是个二进制文件，得用 CUDA Binary Utilities 里的工具打开：

$ cuobjdump -lelf kernel.cubin
ELF file    1: kernel.sm_120.cubin

说明它里面有一个针对 sm_120 的 GPU binary。如果想进一步看看它的 SASS 长什么样（感兴趣的读者可以对照 Blackwell 指令集）：

$ cuobjdump -sass kernel.cubin

        code for sm_120
        .target sm_120

                Function : _Z6VecAddPfS_S_
        .headerflags    @"EF_CUDA_SM120 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM120)"
        /*0000*/                   LDC R1, c[0x0][0x37c] ;                /* 0x0000df00ff017b82 */
                                                                          /* 0x000fe20000000800 */
        /*0010*/                   S2R R9, SR_TID.X ;                     /* 0x0000000000097919 */
                                                                          /* 0x000e2e0000002100 */
        /*0020*/                   LDC.64 R2, c[0x0][0x380] ;             /* 0x0000e000ff027b82 */
                                                                          /* 0x000e220000000a00 */
        /*0030*/                   LDCU.64 UR4, c[0x0][0x358] ;           /* 0x00006b00ff0477ac */
                                                                          /* 0x000e6e0008000a00 */
        /*0040*/                   LDC.64 R4, c[0x0][0x388] ;             /* 0x0000e200ff047b82 */
                                                                          /* 0x000eb00000000a00 */
        /*0050*/                   LDC.64 R6, c[0x0][0x390] ;             /* 0x0000e400ff067b82 */
                                                                          /* 0x000ee20000000a00 */
        /*0060*/                   IMAD.WIDE.U32 R2, R9, 0x4, R2 ;        /* 0x0000000409027825 */
                                                                          /* 0x001fc800078e0002 */
        /*0070*/                   IMAD.WIDE.U32 R4, R9.reuse, 0x4, R4 ;  /* 0x0000000409047825 */
                                                                          /* 0x044fe400078e0004 */
        /*0080*/                   LDG.E R2, desc[UR4][R2.64] ;           /* 0x0000000402027981 */
                                                                          /* 0x002ea8000c1e1900 */
        /*0090*/                   LDG.E R5, desc[UR4][R4.64] ;           /* 0x0000000404057981 */
                                                                          /* 0x000ea2000c1e1900 */
        /*00a0*/                   IMAD.WIDE.U32 R6, R9, 0x4, R6 ;        /* 0x0000000409067825 */
                                                                          /* 0x008fc800078e0006 */
        /*00b0*/                   FADD R9, R2, R5 ;                      /* 0x0000000502097221 */
                                                                          /* 0x004fca0000000000 */
        /*00c0*/                   STG.E desc[UR4][R6.64], R9 ;           /* 0x0000000906007986 */
                                                                          /* 0x000fe2000c101904 */
        /*00d0*/                   EXIT ;                                 /* 0x000000000000794d */
                                                                          /* 0x000fea0003800000 */
        /*00e0*/                   BRA 0xe0;                              /* 0xfffffffc00fc7947 */
                                                                          /* 0x000fc0000383ffff */
        /*00f0*/                   NOP;                                   /* 0x0000000000007918 */
                                                                          /* 0x000fc00000000000 */
                ..........

3.3 symbol level

对于 kern 来说，比较重要的是这个命令：

cuobjdump -symbols kernel.cubin

我们更在意它里面包含的 “GPU Function”，结果是这样的：

symbols:
STT_OBJECT       STB_WEAK   STV_DEFAULT  U .nv.reservedSmem.offset0
STT_?            STB_WEAK   STO_RESERVED_SHARED   __nv_reservedSMEM_offset_0_alias
STT_FUNC         STB_GLOBAL STO_ENTRY      _Z6VecAddPfS_S_

只看最后一行，也就是我们的 VecAdd，各列含义如下：

字段	含义
`STT_FUNC`	函数
`STB_GLOBAL`	全局可见 symbol
`STO_ENTRY`	CUDA kernel entry point
`_Z6VecAddPfS_S_`	kernel 的 ELF symbol name

这个 _Z6VecAddPfS_S_ 非常不人类可读，它是什么东西呢？这是 Itanium C++ ABI 的 name mangling 编码，可以用一个 CLI 解码出它的 “C++ 签名”：

$ cu++filt _Z6VecAddPfS_S_
VecAdd(float *, float *, float *)

这就是后面 kern 的一大根基：我们只需要拿到 GPU binary，就可以反推出很多 kernel 原本的签名，然后执行。

4. 例外：CuTe DSL 的 symbol

当然也有例外。我们用 CuTe DSL（Python）写一个 vector add，然后编译出来：

@cute.kernel
def kernel(a: cute.Tensor, b: cute.Tensor, out: cute.Tensor, n: Int32):
    tidx, _, _ = cute.arch.thread_idx()
    if tidx < n:
        out[tidx] = a[tidx] + b[tidx]

@cute.jit
def vec_add(a: cute.Tensor, b: cute.Tensor, out: cute.Tensor, n: Int32):
    kernel(a, b, out, n).launch(grid=[1, 1, 1], block=[N, 1, 1])

它的 symbol 长这个样子：

STT_FUNC         STB_GLOBAL STO_ENTRY      kernel_cutlass_kernel_tensorptrf32gmemalign16o1_tensorptrf32gmemalign16o1_tensorptrf32gmemalign16o1__0

这个名字是 CuTe DSL 自己定义的编码，不走 Itanium ABI，cu++filt 对它无能为力。大概意思如下：

kernel_                 # CuTe DSL 生成的 CUDA kernel
cutlass_                # CUTLASS DSL mangling prefix
kernel                  # 原始 Python @cute.kernel 的函数名

tensorptr               # cute.Tensor，engine 是 pointer
f32                     # element type = float32
gmem                    # global memory
align16                 # pointer alignment = 16 bytes
o1                      # tensor layout 信息，这里基本是 1D contiguous stride=1

× 3                     # 三个这样的参数

_0                      # CuTe DSL 内部 kernel instance 编号

反解出来大概是这么个东西：

kernel(
    Tensor<f32, gmem, align=16, stride=1>,
    Tensor<f32, gmem, align=16, stride=1>,
    Tensor<f32, gmem, align=16, stride=1>
)

5. dump 一个程序实际用到的 cubin

这里还有一个有意思的话题：如何捕获一个程序执行到的 cubin。比如你拿到了一个很大的 binary，里面有各种五花八门的 kernel，但可能我们只用了其中一部分，显然我们不想把所有东西都 “解码” 出来。再比如我们想拿到 vllm serve Qwen3-4B 实际用到的算子。

目标是：在不改动程序的前提下，拿到这个程序执行过的 cubin。

思路是利用 NV 给 profiler 做的 injection 生态。我们并不只是 “profile”，而是把它执行过的 cubin 都 dump 下来：

在目标 CUDA 程序里插一个观察器
观察 “加载了哪个 cubin” 和 “执行了哪个 kernel”

具体代码比较丑陋（一堆 CUPTI 名词），这里只展示最终效果：

./dump_cubin/dump_cubin --out ./out -- ../kernel

这个 kernel 可以是任意在执行 GPU 的程序。跑下来的日志（加了注释）：

// 准备工作
[dump_cubin] InitializeInjection  (CUDA driver dlsym'd this before the first CUDA call)
[dump_cubin]   CUDA_INJECTION64_PATH=./libdump_cubin.so
[dump_cubin]   DUMP_CUBIN_DIR=./out

// 订阅两类事件：加载 GPU module，launch kernel
[dump_cubin] CUPTI subscribed: RESOURCE/MODULE_LOADED + DRIVER_API/cuLaunchKernel{,Ex}

// 初始化 driver
[dump_cubin] callback  domain=RESOURCE  cbid=CU_INIT_FINISHED
[dump_cubin]   ignore (not a module load)
[dump_cubin] callback  domain=RESOURCE  cbid=CONTEXT_CREATED
[dump_cubin]   ignore (not a module load)

// 调用 cuLaunchKernel
[dump_cubin] callback  domain=DRIVER_API  cbid=cuLaunchKernel  site=ENTER
[dump_cubin]   skip dump: CUDA 12 lazy-load may still be materializing the module
// 开始加载 module
[dump_cubin] callback  domain=RESOURCE  cbid=MODULE_LOADED
[dump_cubin]   runtime just installed a module for this GPU (not a static cuobjdump of the host binary)
[dump_cubin]   moduleId=21  cubinSize=5504  magic=ELF
// 落盘这个 cubin
[dump_cubin]   copied pCubin → ./out/module_21.cubin
[dump_cubin] callback  domain=RESOURCE  cbid=MODULE_PROFILED
[dump_cubin]   ignore (not a module load)
// cuLaunchKernel 正式结束
[dump_cubin] callback  domain=DRIVER_API  cbid=cuLaunchKernel  site=EXIT
// 就是我们之前看到的那个名字
[dump_cubin]   EXIT: launch completed, module is live, symbol=_Z6VecAddPfS_S_

至此 dump 结束，你可以在磁盘上看到一个新的 .cubin 文件了。

值得一提的是 skip dump 那一行：CUDA 12 之后 module 默认 lazy load，cuLaunchKernel 的 ENTER 时刻 module 可能还没 materialize，所以真正的落盘要等 MODULE_LOADED 事件，而不是在 launch 入口抄。

6. 只靠 cubin + symbol 执行：Driver API

铺垫了这么多，我们为什么费劲搞这个 cubin 呢？是为了完全不感知 kernel 实现细节，做更加纯粹的执行器。

有了一个 cubin 和 symbol 后怎么加载它执行？需要 cuModuleLoad，大概流程是这样的：

cuInit(0);

CUdevice dev;
cuDeviceGet(&dev, 0);

CUcontext ctx;
cuCtxCreate(&ctx, 0, dev);

CUmodule module;
cuModuleLoad(&module, "./out/module_21.cubin");

CUfunction func;
cuModuleGetFunction(
    &func,
    module,
    "_Z6VecAddPfS_S_"
);

// 省略一些准备
cuLaunchKernel(
    func,
    1, 1, 1,        // grid
    N, 1, 1,        // block
    0,              // dynamic shared memory
    nullptr,        // stream
    args,           // void*[]，每项指向一个参数
    nullptr
);

注意 args 只是一个 void* 数组，每一项指向一个参数的内存。参数的大小、对齐、偏移都由 kernel image 自身的元数据定义，driver 自己会去读。launch 端只需要按签名把值摆好，不需要知道这个 kernel 是 nvcc、Triton 还是 CuTe DSL 生成的。

具体看 cuModuleLoad，它支持加载这几类文件：

cubin
手写 PTX 文件
fatbin 文件
Tile IR 文件

前三个没什么意外，只是最后这个 Tile IR 是什么东西呢？见 Tile IR bytecode，细节也不展开了，总之就是也可以给一个类似 “PTX” 的东西给 driver 执行。

所以当我们说一个引擎 “支持了 Qwen3 4B”，它大概率只是引入或者复用了一些相关的 kernel，然后把它们的 buffer 连线对接好。仅此而已。

7. kern 的 manifest：引擎只感知四层声明

所以 kern 最初的想法就是：把一个模型的 “图” 描述出来。现在 kern 的一个例子大概是这样的，先是一些常量声明：

{
 "schema_version": 4,
 "model": "qwen3-4b-dspark",
 "vars": {
  "tokens": { "max": 2048 },   // chunk prefill 上限
  "seqs":   { "max": 256 }     // max batch size
 }
}

7.1 states：KV cache 布局

引擎还要感知什么细节呢？这个模型的 KV cache。普通 full attention + draft KV cache 长这样：

"states": {
 "kv":       { "bytes_per_token": 147456 },
 "draft_kv": { "bytes_per_token": 20480 }
}

对于 Qwen3.8 + DFlash2（带 GDN 这种 per-seq 的 recurrent state），长这样：

"states": {
 "kv":       { "bytes_per_token": 65536 },
 "gdn":      { "bytes_per_seq": 163577856 },
 "draft_kv": { "bytes_per_token": 20480 }
}

这是引擎需要感知的 KV cache 细节：按 token 还是按 seq 计量，每单位多少字节。引擎会据此对不同类型的 state 做分配、缓存、offload。

7.2 buffers：一次 forward 的全部中间量

之后是 buffer 声明，即一次 forward 过程中需要用到的全部 buffer，带 dtype、shape 等。这个也可以用来给引擎规划显存：中间激活值完全可以静态推算，扣除少部分 CUDA graph 占用即可。

7.3 modules：kernel 从哪来

然后是声明各种 kernel 的来源。kern 完全不关心这个 kernel 是怎么编写的，不关心具体细节，可以从各种地方获取，包括 Hugging Face 的 kernels 库：

"argmax": {
 "source": "argmax.cubin",
 "sha256": "25374575bd35dbd9fe07f90982d0da0c2d4a72cc6ac88df5ec92cdb62a3beb16"
},
"cache": {
 "source": "cache.cubin",
 "sha256": "93d7768f15ce02e252ad208dca2f6803d8ad65296d57f3444dccb9d1e6b88ac5"
},
"chunk_h": {
 "source": "chunk_h.cubin",
 "sha256": "f0b2bd305bc428b03b06f7073efb3c73d0161fc8c1613c1047e169eae423f32c"
},

// 直接声明 HF 上的路径也是 ok 的
"activation": {
 "source": "hf:kernels-community/activation/build/torch29-cxx11-cu130-aarch64-linux/activation/_activation_320b408.abi3.so",
 "sha256": "73748b54059552f5983322f7dedc36ed349b38ad6fb9318301bb4965b1fe49aa"
}

7.4 ops：kernel 的签名与 launch 参数

光有来源不够，kern 还需要知道这些 kernel 的参数类型：它接受什么类型的输入，用于后续校验 kernel(args) 是否合法。

所以我们需要声明 ops，引用 module，module 里有什么 kernel。比如一个 rms_norm 的声明：

"ops": {
 "rms_norm": {
  "params": [
   "out buffer<bf16>",
   "in buffer<bf16>",
   "in buffer<bf16>",
   "i32"
  ],
  "impl": {
   "launches": [
    {
     "module": "vllm_layernorm",
     "entry": "_ZN4vllm15rms_norm_kernelIN3c108BFloat16ELi8ELi2ELb1EEEvPT_PKS3_lllllS6_lfii",
     "params": [
      "out buffer<bf16>",
      "in buffer<bf16>",
      "i64", "i64", "i64", "i64", "i64",
      "in buffer<bf16>",
      "i64",
      "f32",
      "i32",
      "i32"
     ],
     "block": [320, 1, 1],
     "grid": ["tokens", 1, 1],
     "args": [
      { "param": 0 },
      { "param": 1 },
      { "i64": 2560 },
      { "i64": 0 },
      { "i64": 0 },
      { "i64": 0 },
      { "i64": 0 },
      { "param": 2 },
      { "i64": 0 },
      { "f32": 9.999999974752427e-07 },
      { "param": 3 },
      { "i32": 2560 }
     ]
    }
   ]
  }
 }
}

一个 op 对外是 4 个参数的 rms_norm，对内是一次 launch：指定 module、entry symbol、kernel 真实签名、grid/block、以及每个 kernel 形参从哪里取值（op 的第几个参数，或者一个常量）。这里的 entry 就是第 3 节里那种 mangled symbol，签名也是从它 demangle 出来对上的。

manifest 里就是一大堆这样的 ops 声明。

7.5 program：把 ops 串成链

但这些都是声明，最后怎么执行呢？最后一个抽象就是 program，把这些 ops 拼成一个链条即可。其中 buf 和 var 都是前面提前声明过的：

"calls": [
 {
  "label": "embed",
  "op": "embedding",
  "args": [
   { "buf": "token_ids" },
   { "buf": "model.embed_tokens.weight" },
   { "buf": "residual" },
   { "var": "tokens" }
  ]
 },
 {
  "label": "l0.input_norm",
  "op": "rms_norm",
  "args": [
   { "buf": "x" },
   { "buf": "residual" },
   { "buf": "model.layers.0.input_layernorm.weight" },
   { "var": "tokens" }
  ]
 },
 {
  "label": "l0.qkv_proj",
  "op": "gemm",
  "args": [
   { "buf": "x" },
   { "buf": "model.layers.0.self_attn.qkv_proj.weight" },
   { "buf": "qkv" },
   { "var": "tokens" },
   { "i32": 6144 },
   { "i32": 2560 }
  ]
 }
]

所以最终 kern 如果要 day0 支持一个模型，只要它不在 KV cache 上玩出不一样的花活（比如 v4 flash 那样的不同层级压缩、召回等），理论上 kern 不用改一行代码。

模型提供商只需要公开一个 Hugging Face repo，里面有 weights、model.json、.cubin（可以是 fat 的，也可以不是），kern 就可以立刻开始服务，支持一切你想要的：prefix cache、KV cache offloading。调度也很简单，对 kern 来说就是分配好 buffer，然后推进 program，仅此而已。

8. 除了 day0，kern 还有什么好处

看到这里你又会问了：主播主播，我也不用支持 day0 模型，kern 还有什么好处呢？

快速的算子性能与正确性校验。 比如 rms norm，只要签名没改，算子哥可以不断优化他的 .cu 或者 CuTe DSL，引擎侧不用感知，也不用为他的实现变复杂而改代码。他觉得 ok 以后，就给模型的那个 Hugging Face 仓库提交，更新 cubin 的同时更新 json。

kern 支持 kern test a.json b.json 一键输出三类 diff，帮助算子作者更快完成 verify：

kernel diff：构造不同随机输入分布，逐 op 对比输出
精度 diff：logits 分布是否被扰动
performance diff

调精度也是一把好手。 假设未来生态跟着 kern 走，模型厂商一定有一个 day0 的 json，不管性能怎么样，它的精度至少是对的。内部可能有一次 push 没有做精度校验，积累到某天，服务开始乱码。

这个时候依旧是 diff day0 json 和 optimized json，也是一键 kern test 就好（虽然还没做），停在第一个对不上的 op。

因为 kernel 优化无非几类：

换了个更好的实现，但接口不变。这个很棒，逐 op 对比就够了。
breaking 的融合：A + B → C。这时需要对比的是 AB 的联合输出和 C。
有时候又反过来拆分：C → AB。

不管哪一类，对比对象都是 program 上的一段区间，而 program 本身就是 manifest 里的一个数组。这就是把引擎的感知面压缩到 “四层声明” 之后，顺手得到的东西。