4、CUDA程序编译与执行：nvcc编译器、PTX与SASS、运行时API与驱动API
好，咱们今天聊聊CUDA程序到底是怎么从代码变成实际在GPU上跑的东西。很多人写了好几年CUDA，其实对编译过程还是一知半解。我刚开始做GPU加速的时候也是这样，觉得能跑就行，管它中间经历了什么。直到有一次，我在一个老旧的Tesla K80上调试一个kernel，死活跑不出预期性能，最后发现是编译器生成了不兼容的SASS指令——嗯，从那以后，我再也不敢跳过编译原理了。
4.1 nvcc编译器的工作流程
nvcc，说白了就是NVIDIA的CUDA编译器前端。它不像gcc那样直接一股脑把源码编译成机器码。它很聪明，会把你的.cu文件拆成两部分：
- 主机代码（运行在CPU上）—— 交给宿主编译器（比如gcc、cl.exe）处理
- 设备代码（运行在GPU上）—— 由nvcc自己编译成PTX或SASS
我个人习惯把nvcc的编译过程分成三个阶段：
- 预处理与分离：nvcc先解析.cu文件，把主机代码和设备代码分开。它会生成一个.cu.cpp.ii文件（主机部分）和一个.cu.sm_xx.ptx文件（设备部分）。
- 设备代码编译：设备代码先被编译成PTX（并行线程执行）中间表示。PTX是一种虚拟指令集，不依赖具体GPU架构。
- PTX到SASS：在运行时，PTX会被进一步编译成SASS（流式汇编），这才是GPU真正执行的机器码。这一步由显卡驱动中的JIT编译器完成。
核心要点：nvcc的编译过程是两阶段的。第一阶段生成PTX，第二阶段在运行时将PTX编译为SASS。这种设计保证了CUDA程序的向前兼容性。
4.2 PTX与SASS：中间表示与机器码
PTX和SASS，这两个概念我经常看到有人搞混。你想想看，PTX就像Java的字节码，而SASS就是真正的x86机器码。PTX是跨架构的，SASS是架构相关的。
| 特性 | PTX | SASS |
|---|---|---|
| 生成时机 | 编译时（nvcc -ptx） | 运行时（JIT编译） |
| 架构依赖 | 不依赖具体架构 | 依赖具体架构（如sm_75） |
| 可读性 | 较好，类似汇编 | 较差，二进制格式 |
| 优化级别 | 中等优化 | 深度优化（寄存器分配、指令调度） |
| 用途 | 跨平台分发、调试 | 实际执行 |
我在项目中遇到过一个问题：用nvcc编译时指定了-arch=sm_75，但部署的机器是RTX 3090（sm_86）。按理说应该报错，但程序居然能跑。为什么？因为nvcc默认会嵌入PTX代码，驱动发现SASS不兼容，就自动用PTX重新JIT编译了。这就是PTX的妙处。
实用技巧：编译时建议同时指定-arch和-code。比如：nvcc -arch=sm_75 -code=sm_75,sm_86。这样既生成了sm_75的SASS，也嵌入了PTX，保证向后兼容。
4.3 运行时API与驱动API
CUDA提供了两套API来跟GPU打交道：运行时API和驱动API。我刚开始学CUDA时，用的全是运行时API，因为它简单。后来做底层优化时，才发现驱动API的威力。
运行时API（Runtime API）
这是最常用的API。你写的cudaMalloc、cudaMemcpy、cudaLaunchKernel都属于运行时API。它封装了驱动API，使用起来更简单。说白了，运行时API就是驱动API的“保姆版”。
// 运行时API示例
cudaMalloc(&d_data, size);
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
myKernel<<<grid, block>>>(d_data);
cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);
cudaFree(d_data);驱动API（Driver API）
驱动API更底层，直接跟CUDA驱动打交道。它需要你手动管理上下文、模块、函数句柄。听起来麻烦，但好处是控制力更强。我记得有一次做多GPU编程，需要精确控制每个GPU的上下文切换，运行时API根本做不到，只能用驱动API。
// 驱动API示例
CUdevice device;
CUcontext context;
CUmodule module;
CUfunction kernel;
cuDeviceGet(&device, 0);
cuCtxCreate(&context, 0, device);
cuModuleLoad(&module, "mykernel.ptx");
cuModuleGetFunction(&kernel, module, "myKernel");
void *args[] = { &d_data };
cuLaunchKernel(kernel, grid.x, grid.y, grid.z,
block.x, block.y, block.z,
0, NULL, args, NULL);避坑指南：我曾经在项目中混合使用运行时API和驱动API，结果导致上下文冲突，程序崩溃。记住：一个进程中，运行时API和驱动API不能混用上下文。要么全用运行时，要么全用驱动。
4.4 编译选项与优化技巧
nvcc的编译选项很多，我挑几个最常用的说说：
- -arch：指定虚拟架构，比如
-arch=sm_75。决定了生成的PTX指令集版本。 - -code：指定实际架构，比如
-code=sm_86。决定了生成的SASS目标。 - -G：生成调试信息，但会关闭优化。调试时用，发布时千万别加。
- -O3：最高优化级别。我一般默认加这个。
- --use_fast_math：使用快速数学函数，牺牲精度换速度。适合对精度要求不高的场景。
我的建议：发布程序时，最好同时编译多个架构的SASS，并嵌入PTX。这样既能保证性能，又能兼容新旧显卡。命令示例：
nvcc -arch=sm_75 -code=sm_75,sm_86,sm_89 -O3 --use_fast_math mykernel.cu -o mykernel4.5 实战：查看PTX和SASS
有时候你需要看看编译器到底生成了什么代码。我常用的方法：
- 查看PTX：
nvcc -ptx mykernel.cu，会生成mykernel.ptx文件。 - 查看SASS：
nvcc -cubin -arch=sm_75 mykernel.cu，然后用cuobjdump -sass mykernel.cubin反汇编。
你想想看，当你发现kernel性能不对劲时，看看生成的SASS，往往能发现编译器做了哪些优化（或者没做哪些优化）。比如，我上次发现一个循环没有被完全展开，就是因为SASS里看到了多余的跳转指令。
调试技巧：用cuobjdump -ptx mykernel可以查看可执行文件中嵌入的PTX代码。这在你拿到别人的二进制文件时特别有用。
好了，这一章的内容就到这里。记住：理解编译过程，能帮你写出更高效的CUDA程序，也能帮你更快地定位问题。下一章我们会聊聊GPU的内存层次结构，那才是真正决定性能的关键。