nvcc 是 NVIDIA CUDA 生态的核心编译器，负责将 CUDA C/C++ 代码（混合了主机代码和设备代码）编译为可在 CPU 和 GPU 上运行的二进制文件。它不仅是一个简单的编译器，更是一个“编译驱动程序”，协调多个工具链（如主机编译器、CUDA 设备编译器、汇编器、链接器）完成整个编译流程。

一、nvcc 核心功能与编译流程

1. 核心功能

• 分离并处理 主机代码（CPU 执行，如 main 函数）和 设备代码（GPU 执行，如 __global__ 函数）。
• 将设备代码编译为 GPU 可执行的二进制指令（SASS）或中间代码（PTX）。
• 协调主机编译器（如 gcc、cl.exe）处理主机代码，并链接 CPU/GPU 代码生成最终可执行文件。

2. 编译流程（简化版）

源代码（.cu）↓ 预处理（合并头文件、展开宏等）
预处理文件（.i）↓ 前端编译（分离主机/设备代码）├─ 主机代码 → 交给主机编译器（如 gcc）编译为 CPU 目标文件└─ 设备代码 → 编译为 PTX 中间代码（.ptx）↓ 后端编译（PTX → GPU 二进制指令 SASS）设备目标文件（.o）↓ 链接（合并主机/设备目标文件 + 链接 CUDA 库）
最终可执行文件

二、基础用法与核心选项

1. 基础编译命令

# 直接编译 .cu 文件为可执行文件
nvcc source.cu -o output# 分步编译（先编译为目标文件，再链接）
nvcc -c source.cu -o source.o  # 生成目标文件
nvcc source.o -o output        # 链接生成可执行文件

2. 架构相关选项（最核心！）

CUDA 代码需针对特定 GPU 架构优化，核心选项如下：

选项	作用	示例
-arch=sm_xx	指定目标 GPU 的计算能力（生成该架构的 PTX 和 SASS）	-arch=sm_75（图灵架构）
-code=sm_xx	仅生成指定架构的 SASS 代码（需与 -arch 配合）	-arch=compute_70 -code=sm_70,sm_75
-gencode=...	生成多种架构的代码（更灵活的多架构支持）	-gencode arch=compute_70,code=sm_70
--list-gpu-arch	查看当前 CUDA 版本支持的所有架构	nvcc --list-gpu-arch

示例：编译支持多架构的程序
（同时支持图灵 sm_75 和安培 sm_86）：

nvcc -gencode arch=compute_75,code=sm_75 \-gencode arch=compute_86,code=sm_86 \source.cu -o output

注意：compute_xx 表示虚拟架构（用于生成 PTX），sm_xx 表示实际架构（用于生成 SASS）。

3. 预处理选项

控制预处理阶段（类似 gcc 的 -E、-D 等）：

选项	作用	示例
-E	仅执行预处理，输出预处理后的代码	nvcc -E source.cu -o source.i
-D<宏名>	定义宏（编译时生效）	-DDEBUG（启用调试宏）
-I<路径>	添加头文件搜索路径	-I./include
-U<宏名>	取消已定义的宏	-UDEBUG

4. 编译与优化选项

控制编译阶段的优化、调试等行为：

选项	作用	示例
-O0~-O3	优化级别（-O0 无优化，-O3 最高优化，默认 -O1）	-O2
-g	保留调试信息（供 gdb 等工具使用）	-g
-G	生成设备代码调试信息（会关闭部分优化，仅用于调试）	-G
-lineinfo	保留设备代码行号信息（供 nvprof 等性能分析工具使用，不影响优化）	-lineinfo
--use_fast_math	启用快速数学库（精度略低，速度更快）	--use_fast_math
-Xptxas <选项>	向 PTX 到 SASS 的汇编器传递选项（如控制寄存器使用）	-Xptxas -v（输出汇编详细信息）

5. 链接选项

控制链接阶段的库依赖和路径：

选项	作用	示例
-L<路径>	添加库文件搜索路径	-L/usr/local/cuda/lib64
-l<库名>	链接指定库（如 CUDA 运行时、cuBLAS 等）	-lcudart（链接 CUDA 运行时库）
-shared	生成动态链接库（.so/.dll）	nvcc -shared -fPIC source.cu -o libxxx.so
-static	静态链接（仅主机代码，设备代码无法静态链接）	-static

6. 输出中间文件选项

生成编译过程中的中间文件（用于调试或分析）：

选项	作用	示例
-ptx	生成 PTX 中间代码（GPU 汇编的文本形式）	nvcc -ptx source.cu（输出 source.ptx）
-cubin	生成特定架构的二进制代码（.cubin）	nvcc -arch=sm_75 -cubin source.cu
-keep	保留所有中间文件（.i、.ptx、.cubin 等）	nvcc -keep source.cu

三、进阶场景与示例

1. 混合主机代码（C++）与设备代码（CUDA）

假设项目结构：

project/
├── main.cpp    # 主机代码（C++）
└── kernel.cu   # 设备代码（CUDA）

编译步骤：

# 1. 编译设备代码为目标文件（.o）
nvcc -arch=sm_75 -c kernel.cu -o kernel.o# 2. 用 g++ 编译主机代码（需包含 CUDA 头文件）
g++ -c main.cpp -o main.o -I/usr/local/cuda/include# 3. 链接（用 nvcc 确保 CUDA 库正确链接）
nvcc kernel.o main.o -o mixed_app

2. 生成动态链接库（.so）

将 CUDA 代码封装为可被其他程序调用的动态库：

# 编译为动态库（-fPIC 生成位置无关代码）
nvcc -arch=sm_75 -shared -fPIC cuda_lib.cu -o libcudax.so# 使用该库（例如在 C++ 程序中调用）
g++ app.cpp -o app -L. -lcudax -lcudart  # -lcudart 链接 CUDA 运行时

3. 交叉编译（例如在 x86 主机编译 ARM 目标）

需配合 NVIDIA 交叉编译工具链（如 Jetson 平台）：

# 假设交叉编译器为 aarch64-linux-gnu-g++
nvcc -arch=sm_72 \-ccbin aarch64-linux-gnu-g++ \  # 指定主机交叉编译器-target-cpu-arch aarch64 \      # 目标 CPU 架构source.cu -o output

4. 查看详细编译过程（调试编译问题）

使用 -v 选项输出编译全过程（包括调用的工具、参数等）：

nvcc -v source.cu -o output  # 详细输出编译步骤

四、与构建工具结合（Makefile/CMake）

1. 示例 Makefile

NVCC = nvcc
ARCH = -arch=sm_75
OPTFLAGS = -O2 -lineinfo
INCLUDES = -I./include
LIBS = -lcublas -lcufft  # 链接 cuBLAS 和 cuFFT 库SRC = main.cu kernel.cu
OBJ = $(SRC:.cu=.o)
TARGET = cuda_appall: $(TARGET)%.o: %.cu$(NVCC) $(ARCH) $(OPTFLAGS) $(INCLUDES) -c $< -o $@$(TARGET): $(OBJ)$(NVCC) $(OBJ) -o $@ $(LIBS)clean:rm -f $(OBJ) $(TARGET)

2. 示例 CMakeLists.txt

cmake_minimum_required(VERSION 3.18)
project(cuda_example)# 寻找 CUDA 工具包
find_package(CUDA 12.0 REQUIRED)# 设置 GPU 架构
set(CUDA_ARCHITECTURES 75 86)  # 支持 sm_75 和 sm_86# 添加可执行目标
cuda_add_executable(cuda_app main.cu kernel.cu)# 链接 CUDA 库（如 cuBLAS）
target_link_libraries(cuda_app CUDA::cublas)

五、常见问题与调试技巧

1. “nvcc: command not found”

• 原因：CUDA 未安装或环境变量未配置。
• 解决：添加环境变量（以 CUDA 12.2 为例）：

  export PATH=/usr/local/cuda-12.2/bin:$PATH
  export LD_LIBRARY_PATH=/usr/local/cuda-12.2/lib64:$LD_LIBRARY_PATH
  

2. “unsupported gpu architecture ‘sm_xx’”

• 原因：指定的 sm_xx 不在当前 CUDA 版本支持范围（如 CUDA 12 不支持 sm_20）。
• 解决：用 nvcc --list-gpu-arch 查看支持的架构，更换为兼容值。

3. 设备函数未定义（“undefined reference to `kernel<<<>>>'”）

• 原因：设备代码未被正确编译或链接（如用 g++ 直接链接 .cu 文件）。
• 解决：确保用 nvcc 编译 .cu 文件，且链接时使用 nvcc 而非主机编译器。

4. 优化效果不佳

• 技巧：
  • 用 -Xptxas -v 查看寄存器使用情况（寄存器不足会导致性能下降）。
  • 启用 -O3 和 --use_fast_math（根据精度需求）。
  • 配合 nvprof 或 nsys 分析性能瓶颈：

    nvprof ./output  # 基础性能分析
    nsys profile ./output  # 更详细的全系统分析
    

六、总结

nvcc 的核心是“协调主机与设备代码的编译流程”，掌握它的关键在于：

1. 正确指定 GPU 架构（-arch/-gencode），匹配目标硬件。
2. 合理使用优化选项（-O2/-Xptxas）提升性能。
3. 结合构建工具（Makefile/CMake）管理复杂项目。
4. 利用中间文件（PTX/SASS）和调试选项（-v/-G）解决编译问题。

通过 nvcc --help 可查看完整选项列表，结合实际项目调试能更快掌握其用法。