【CUDA】cubin是什么

简略

cubin = CUDA binary。

在 CUDA/NVCC 语境里，它通常表示：针对某个 GPU 架构已经编译好的设备代码二进制。

cubin vs <<<>>>

cubin：

kernel 函数不作为“可直接用 <<<>>> 调用的代码”放在主程序里，而是单独生成一个设备二进制模块文件；运行时先 load 模块，再按名字取 kernel，再 launch。可类比动态库。

非 cubin（<<<>>> 路径）：

kernel 和 host 一起编译，设备代码打包进最终程序；运行时直接写 <<<>>>，由 runtime/driver 在背后完成内核定位、装载和 launch。可以类比静态库。

1. 设备代码二进制放在哪里

• <<<>>>：通常内嵌在最终程序里
• cubin：单独放成外部模块文件

2. 运行时由谁加载、怎么启动

• <<<>>>：runtime/driver 自动处理
• cubin：程序显式 load module + get function + launch

详细说明

1）<<<>>> 路径

比如 nvshmem 的 put_signal.cu:109-110：

test_##FUNC##_signal##SC_SUFFIX##_kernel<<<1, num_threads>>>(
    target, source, sig_addr, me, npes, NVSHMEM_SIGNAL_SET);

• kernel 源码和 host 代码一起编译
• 编译器把 device 端二进制 打包进最终程序里
• 运行时直接写 <<<>>>
• runtime/driver 在背后帮你完成：

• • 找到这个 kernel
  • 装载对应设备代码
  • 发起 launch

2）“cubin 路径”

比如这个文件里：

• put_signal.cu:12 定义了 put_signal.cubin
• put_signal.cu:164-166 里 init_cumodule(CUMODULE_NAME)
• put_signal.cu:24-31 最后 cuLaunchKernel(...)

这条路的意思是：

• device 代码先单独编译成一个外部模块文件
• 运行时显式 load module
• 再按名字取出 kernel
• 再显式 cuLaunchKernel

所以它的本质是：

设备二进制独立放在外部模块里，由程序自己显式加载并启动

2. cubin 和 <<<>>> 的区别

方式 A：原生 <<<>>>

这是最常见的 CUDA/HIP 风格：

kernel<<<grid, block>>>(args);

特点：

• 源码里直接写 kernel 名字
• 编译时就把 kernel 和主程序一起处理好了
• 运行时不用你自己手动加载 module
• 用起来最直接

在这个测例里就是这条路径 (put_signal.cu:109-110)：

test_##FUNC##_signal##SC_SUFFIX##_kernel<<<1, num_threads>>>(
    target, source, sig_addr, me, npes, NVSHMEM_SIGNAL_SET);

方式 B：cubin + Driver API

这是“先加载模块，再按名字找 kernel，再启动”的方式。

在这个测例里对应的是：

• init_cumodule(CUMODULE_NAME) (put_signal.cu:164-166)
• init_test_case_kernel(...)
• cuLaunchKernel(...) (put_signal.cu:24-31)

也就是类似：

cuModuleLoad(...)
cuModuleGetFunction(...)
cuLaunchKernel(...)

特点：

• kernel 在外部模块文件里
• 运行时显式加载
• 通过字符串名字找函数
• 更接近 driver 层
• 常用于测试“预编译设备代码模块”这条链路是否正常

3. 两条路径调用示例

cubin 路径

put_signal.cu:105-107

if (use_cubin) {
    int op = NVSHMEM_SIGNAL_SET;
    TEST_NVSHMEM_PUT_CUBIN(FUNC, SC_SUFFIX);
}

原生路径

put_signal.cu:108-111

else {
    test_##FUNC##_signal##SC_SUFFIX##_kernel<<<1, num_threads>>>(
        target, source, sig_addr, me, npes, NVSHMEM_SIGNAL_SET);
}

5. 为什么测例里要同时保留这两种？

通常有几个目的：

1）测试 device code module 这条链路

比如：

• 预编译的 .cubin
• 动态加载
• cuLaunchKernel

这条链路在某些场景下是单独需要验证的。

2）覆盖不同编译/链接模式

有些功能在：

• 普通源码直编
• device library / bitcode / cubin 模式

行为可能不一样，所以测试要两边都跑。

3）隔离编译器行为

有时 driver API 路径和 <<<>>> 路径在：

• 资源分配
• 符号解析
• launch 限制
• 编译选项继承

上会有差异，所以测试里会专门比较。

7. 总结

• <<<>>>：直接启动
• cubin：先加载模块，再 cuLaunchKernel

它们是同一类 kernel 的两种装载/启动路径。

类比：

• <<<>>> 写法Kernel 代码和宿主代码一起编译打包，设备二进制内嵌进主程序，运行时自动调用 → 类比 静态库（编译时链接进程序，无需额外加载）。

• Cubin 写法Kernel 单独编译成独立二进制文件，运行时手动加载模块、查找函数再执行 → 类比 动态库（外部文件，运行时显式加载使用）。