nvcc

CUDA 编译器,把 .cu 文件编成能在 GPU 上跑的可执行。每次写 kernel 都绕不开。

下面用的 hello.cu

后面的命令都假设当前目录有 hello.cu,直接复制下面这段保存就行:

// hello.cu
#include <stdio.h>

__global__ void hello() {
    printf("Hello, CUDA from thread %d\n", threadIdx.x);
}

int main() {
    hello<<<1, 8>>>();
    cudaDeviceSynchronize();
    return 0;
}

验证:

nvcc hello.cu -o hello
./hello
# 期望输出 8 行 "Hello, CUDA from thread N"(顺序不固定)

基础编译

nvcc hello.cu -o hello              # 默认架构编译,生成可执行
nvcc -O3 hello.cu -o hello          # 开优化
nvcc -g -G hello.cu -o hello        # 带 host + device 调试信息(给 cuda-gdb 用)

指定 GPU 架构

每张 GPU 有自己的 compute capability(架构号)。不指定的话 nvcc 会用默认值,可能没用到新指令,甚至跑不起来。

nvcc -arch=sm_75 hello.cu -o hello  # Turing(T4 / RTX 20)
nvcc -arch=sm_80 hello.cu -o hello  # Ampere(A100)
nvcc -arch=sm_86 hello.cu -o hello  # Ampere(RTX 30)
nvcc -arch=sm_89 hello.cu -o hello  # Ada Lovelace(RTX 40)
nvcc -arch=sm_90 hello.cu -o hello  # Hopper(H100)

不确定自己显卡是什么架构,先查一下:

nvidia-smi --query-gpu=compute_cap --format=csv

多架构编译(发布常用)

让同一个二进制能在多种卡上跑——开发时不常用,发包给别人时用:

nvcc -gencode arch=compute_75,code=sm_75 \
     -gencode arch=compute_80,code=sm_80 \
     -gencode arch=compute_86,code=sm_86 \
     hello.cu -o hello

看 PTX / SASS

nvcc -ptx hello.cu -o hello.ptx     # 生成中间表示 PTX(类汇编)
nvcc -cubin hello.cu -o hello.cubin # 生成 GPU 二进制 cubin
cuobjdump --dump-sass hello         # 反汇编看真正运行的 SASS 指令

链接 CUDA 库

nvcc main.cu -lcublas -lcudart -o main   # 链 cuBLAS
nvcc main.cu -lcudnn -o main             # 链 cuDNN
nvcc main.cu -lcurand -o main            # 链 cuRAND(GPU 随机数)

注:-arch=sm_xx 是 -gencode arch=compute_xx,code=sm_xx 的简写。前者只产 SASS(机器码),后者产 PTX + SASS,跨架构兼容性更好——发布用 -gencode,自己开发用 -arch= 更简短。