NVCC 编译器:从源码到 GPU 指令

NVCC 是什么?说白了它就是 NVIDIA 的 CUDA 编译器。它的工作就是把你的 .cu 文件,一步步变成 GPU 能直接执行的机器码。我刚开始用 CUDA 时,总觉得编译就是敲个 nvcc -o test test.cu 就完事了。后来踩了几次坑才发现,里面的门道多着呢。

这一节,咱们就把它掰开揉碎了讲清楚。我会从编译流程、关键编译选项,再到 PTX 和 SASS 的区别,一个一个过。

NVCC 的编译流程

一个 .cu 文件要变成 GPU 上的可执行代码,得经过四个阶段。我习惯画一张图记在脑子里,你们也可以试试看。

NVCC 编译流程 预处理 处理宏、头文件 编译 → PTX -arch 控制 PTX 版本 汇编 → SASS -code 指定 GPU 架构 链接 生成 .exe / .so .cu 源文件 .ptx 中间文件 .cubin / 胖二进制 最终可执行 💡 关键概念:虚拟架构 vs 真实架构 -arch 指定的是虚拟架构(例如 sm_50),决定 PTX 能用的指令集。 -code 指定的是真实架构(例如 sm_52),决定最终 SASS 的目标 GPU。 PTX → 统一中间表示(跨代兼容) SASS → 特定GPU机器码(极致性能)

这张图里每个阶段都有讲究。先过一遍流程:

  1. 预处理:处理 #include、宏定义这些脏活。跟 C 编译器的预处理差不多,没啥特殊的。
  2. 编译:把 .cu 里的 __global__ 和 __device__ 函数编译成 PTX 中间表示。这里 -arch 选项决定了输出的 PTX 版本。
  3. 汇编:PTX 被转换成当前 GPU 架构的机器码——SASS。这个过程叫 ptxas,-code 选项告诉它目标 GPU 是哪个。
  4. 链接:把所有 .o 或胖二进制文件链接成最终的可执行文件或库。

我的经验:有时候编译报错说“sm_50 不支持某个 intrinsic”,那就是你 -arch 设低了。我一般会设成项目能支持的最低架构,然后用 -code 列上所有目标 GPU,这样兼容性和性能都能照顾到。

关键的编译选项

NVCC 的选项真的不少。但常用的就那几个。我个人觉得,搞明白 -arch 和 -code 的区别,就解决了 80% 的问题。

1. -arch 与 -code:虚拟与真实的区别

很多人写 -arch=sm_70,然后就不管了。其实这样默认 -code 也等于 sm_70,编译出来的二进制只能在 Volta 及以上架构跑。要是用户用个老显卡,直接崩。

正确的用法是这样:

nvcc -arch=compute_50 -code=sm_50,sm_70,sm_80 test.cu
选项 含义 典型值
-arch 虚拟架构,指定 PTX 指令集版本 compute_30, compute_50, compute_70
-code 真实架构,指定目标 GPU 的 SASS sm_30, sm_50, sm_70
-gencode 灵活组合多组 arch/code arch=compute_50,code=sm_50;arch=compute_70,code=sm_70

一句话总结:-arch 管的是“编译器能用哪些指令”,-code 管的是“最后生成的机器码给哪个 GPU 用”。

2. -G:调试模式

-G 选项会在 PTX 和 SASS 中嵌入大量调试信息。你想想看,如果没有 -G,你在 Nsight 里断点都打不了,变量也看不了。

但我得提醒一句:-G 会让性能大幅下降。我在项目中遇到过这种坑——调试模式跑得好好的,去掉 -G 后结果变了。这是因为 -G 禁用了很多编译器优化。所以记得:调试用 -G,发布一定要去掉。

# 调试版
nvcc -G -o test test.cu

# 发布版(默认优化)
nvcc -O2 -o test test.cu

我曾经在调一个 kernel 的时候,发现加上 -G 后 bug 消失了,去掉又出现了。折腾了两天才发现是 -G 改写了指令顺序,碰巧掩盖了一个未初始化的变量。所以 -G 不是万能的,它只是帮你更方便地找问题。

3. 其他常用选项

  • -O0, -O1, -O2, -O3:控制优化级别。默认是 -O2,适合大多数情况。
  • --ptxas-options=-v:让 ptxas 输出寄存器用量、shared memory 使用量等统计。写优化必加。
  • -keep:保留编译过程中产生的 .ptx、.cubin 等中间文件。适合想看编译器干了些啥。
  • -std=c++17:指定 C++ 标准。CUDA 12 对 C++17 支持很好。

PTX 与 SASS:中间表示与机器码

这两个词经常被搞混。简单说:

  • PTX(Parallel Thread eXecution)是一种伪汇编语言。它不针对特定 GPU,而是定义了一套虚拟指令集。不同代 GPU 都支持同一版本的 PTX,但底层的实现不同。
  • SASS(Streaming Assembly)是真正的 GPU 机器码。每个 GPU 架构(如 Maxwell、Pascal、Turing)都有自己的 SASS 格式。

为什么要搞两层?说白了,为了兼容性。你用 -arch=compute_70 编译出 PTX,可以挂在将来还不存在的架构上,NVCC 在运行时将 PTX JIT 编译成新架构的 SASS。这对软件分发特别友好。

但代价是什么?性能损失。JIT 编译不可能像离线编译那样疯狂优化。所以追求极致性能时,一定要给 -code 指定你真正要跑的 GPU 架构。

我的建议:如果是公司内部部署,硬件统一,那就只用 -arch 和 -code 指定该架构,不要额外塞 PTX。如果是要发布给用户,最好同时包含 PTX 和主流几个 sm_ 的 SASS,让 NVCC 胖二进制自己选。

实践场景:到底该用哪个选项组合?

我来给你几个典型场景的例子,照着套就行:

场景 推荐命令 理由
只跑在 RTX 3090 (Ampere) -arch=compute_80 -code=sm_80 性能最大化,不兼容老卡
兼容 Tesla T4,同时想跑在 V100 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_75,code=sm_75 同时包含两个架构的 SASS
发布给用户,目标未知 -arch=compute_60 -code=sm_60,sm_70,sm_80 -code=compute_80 包含 PTX 兼容未来架构,同时提供三个主流 SASS
纯调试(不关心性能) -G -arch=compute_70 -code=sm_70 调试信息全开,指定一个架构即可

嗯,到这里,NVCC 编译这块就讲完了。你想想看,是不是原来觉得“不就是个编译器”的东西,拆开来看其实挺有深度的?特别是 PTX 和 SASS 的设计哲学,理解了它,你就能更好地平衡兼容性和性能。以后用 NVCC 的时候,选 option 就不会手抖了。