2.7. NVCC：NVIDIA CUDA 编译器

NVIDIA CUDA 编译器nvcc是 NVIDIA 提供的工具链，用于编译 CUDA C/C++ 以及PTX代码。该工具链是CUDA Toolkit的一部分，由多个工具组成，包括编译器、链接器以及 PTX 和Cubin汇编器。顶层工具nvcc负责协调编译过程，为每个编译阶段调用相应的工具。

nvcc驱动 CUDA 代码的离线编译，与之相对的是由 CUDA 运行时编译器nvrtc驱动的在线编译或即时（JIT）编译。

本章介绍构建应用程序所需的nvcc最常见的用法和细节。有关nvcc的完整介绍，请参阅nvcc 文档。

2.7.1. CUDA 源文件和头文件

使用nvcc编译的源文件可以同时包含在 CPU 上执行的主机代码和在 GPU 上执行的设备代码。对于仅包含主机代码的文件，nvcc接受常见的 C/C++ 源文件扩展名.c、.cpp、.cc、.cxx；对于包含设备代码或混合主机代码和设备代码的文件，则接受.cu。包含设备代码的头文件通常采用.cuh扩展名，以区别于仅包含主机代码的头文件.h、.hpp、.hh、.hxx等。

文件扩展名	说明	内容
`.c`	C 源文件	仅主机代码
`.cpp`, `.cc`, `.cxx`	C++ 源文件	仅主机代码
`.h`, `.hpp`, `.hh`, `.hxx`	C/C++ 头文件	设备代码、主机代码、主机/设备混合代码
`.cu`	CUDA 源文件	设备代码、主机代码、主机/设备混合代码
`.cuh`	CUDA 头文件	设备代码、主机代码、主机/设备混合代码

2.7.2. NVCC 编译工作流

在初始阶段，nvcc将设备代码与主机代码分离，并分别将它们的编译工作交给 GPU 编译器和主机编译器。

为了编译主机代码，CUDA 编译器nvcc要求有可用的兼容主机编译器。CUDA Toolkit 定义了Linux和Windows平台的主机编译器支持策略。

仅包含主机代码的文件可以使用nvcc构建，也可以直接使用主机编译器构建。生成的目标文件可以在链接时与nvcc生成的包含 GPU 代码的目标文件组合。

GPU 编译器将 C/C++ 设备代码编译为 PTX 汇编代码。对于编译命令行中指定的每一种虚拟机指令集架构（例如compute_90），都会运行一次 GPU 编译器。
各份 PTX 代码随后被传递给ptxas工具，该工具为目标硬件 ISA 生成Cubin。硬件 ISA 由其SM 版本标识。
可以将多个 PTX 和 Cubin 目标嵌入到应用程序或库中的单个二进制Fatbin容器中，使单个二进制文件能够支持多种虚拟 ISA 和目标硬件 ISA。

上述工具的调用和协调由nvcc自动完成。-v选项可用于显示完整的编译工作流和工具调用。-keep选项可用于将编译过程中生成的中间文件保存在当前目录中，或者改为保存在--keep-dir指定的目录中。

以下示例说明了 CUDA 源文件example.cu的编译工作流：

// ----- example.cu -----
#include <stdio.h>
__global__ void kernel() {
    printf("Hello from kernel\n");
}

void kernel_launcher() {
    kernel<<<1, 1>>>();
    cudaDeviceSynchronize();
}

int main() {
    kernel_launcher();
    return 0;
}

nvcc的基本编译工作流：

nvcc针对多种 PTX 和 Cubin 架构的编译工作流：

有关nvcc编译工作流的更详细说明，请参阅编译器文档。

2.7.3. NVCC 基本用法

使用nvcc编译 CUDA 源文件的基本命令为：

nvcc <source_file>.cu -o <output_file>

nvcc接受常见的编译器标志，用于指定包含目录-I <path>和库路径-L <path>、链接其他库-l<library>以及定义宏-D<macro>=<value>。

nvcc example.cu -I path_to_include/ -L path_to_library/ -lcublas -o <output_file>

2.7.3.1. NVCC 的 PTX 和 Cubin 生成

默认情况下，nvcc为 CUDA Toolkit 支持的最早 GPU 架构（最低的compute_XY和sm_XY版本）生成 PTX 和 Cubin，以最大限度地提高兼容性。

-arch选项可用于为特定的 GPU 架构生成 PTX 和 Cubin。
-gencode选项可用于为多种 GPU 架构生成 PTX 和 Cubin。

要获取所支持的虚拟 GPU 架构和真实 GPU 架构的完整列表，可以分别传递--list-gpu-code和--list-gpu-arch标志，或参阅nvcc文档中的虚拟架构列表和GPU 架构列表部分。

nvcc --list-gpu-code # list all supported real GPU architectures
nvcc --list-gpu-arch # list all supported virtual GPU architectures

nvcc example.cu -arch=compute_<XY> # e.g. -arch=compute_80 for NVIDIA Ampere GPUs and later
                                   # PTX-only, GPU forward compatible

nvcc example.cu -arch=sm_<XY>      # e.g. -arch=sm_80 for NVIDIA Ampere GPUs and later
                                   # PTX and Cubin, GPU forward compatible

nvcc example.cu -arch=native       # automatically detects and generates Cubin for the current GPU
                                   # no PTX, no GPU forward compatibility

nvcc example.cu -arch=all          # generate Cubin for all supported GPU architectures
                                   # also includes the latest PTX for GPU forward compatibility

nvcc example.cu -arch=all-major    # generate Cubin for all major supported GPU architectures, e.g. sm_80, sm_90,
                                   # also includes the latest PTX for GPU forward compatibility

更高级的用法允许分别指定 PTX 和 Cubin 目标：

# generate PTX for virtual architecture compute_80 and compile it to Cubin for real architecture sm_86, keep compute_80 PTX
nvcc example.cu -arch=compute_80 -gpu-code=sm_86,compute_80 # (PTX and Cubin)

# generate PTX for virtual architecture compute_80 and compile it to Cubin for real architecture sm_86, sm_89
nvcc example.cu -arch=compute_80 -gpu-code=sm_86,sm_89    # (no PTX)
nvcc example.cu -gencode=arch=compute_80,code=sm_86,sm_89 # same as above

# (1) generate PTX for virtual architecture compute_80 and compile it to Cubin for real architecture sm_86, sm_89
# (2) generate PTX for virtual architecture compute_90 and compile it to Cubin for real architecture sm_90
nvcc example.cu -gencode=arch=compute_80,code=sm_86,sm_89 -gencode=arch=compute_90,code=sm_90

有关用于控制 GPU 代码生成的nvcc命令行选项的完整参考资料，请参阅nvcc 文档。

2.7.3.2. 主机代码编译注意事项

不包含设备代码或符号的编译单元（即一个源文件及其头文件）可以直接使用主机编译器编译。如果任何编译单元使用 CUDA 运行时 API 函数，则应用程序必须链接 CUDA 运行时库。CUDA 运行时同时提供静态库和共享库，分别为libcudart_static和libcudart。默认情况下，nvcc链接静态 CUDA 运行时库。要使用 CUDA 运行时的共享库版本，请在编译或链接命令中向nvcc传递--cudart=shared标志。

nvcc允许通过-ccbin <compiler>参数指定用于主机函数的主机编译器。也可以定义环境变量NVCC_CCBIN来指定nvcc使用的主机编译器。nvcc的-Xcompiler参数会将参数透传给主机编译器。例如，在下面的示例中，-O3参数由nvcc传递给主机编译器。

nvcc example.cu -ccbin=clang++

export NVCC_CCBIN='gcc'
nvcc example.cu -Xcompiler=-O3

2.7.3.3. GPU 代码的分离编译

nvcc默认采用全程序编译，这种编译方式要求所有 GPU 代码和符号都存在于使用它们的编译单元中。CUDA 设备函数可以调用其他编译单元中定义的设备函数或访问其中定义的设备变量，但必须在nvcc命令行中指定-rdc=true标志或其别名-dc标志，才能启用来自不同编译单元的设备代码的链接。链接不同编译单元中的设备代码和符号的能力称为分离编译。

分离编译允许更灵活地组织代码，可以缩短编译时间，并且可以生成更小的二进制文件。与全程序编译相比，分离编译可能会带来一些构建时的复杂性。使用设备代码链接可能会影响性能，因此默认情况下不使用它。链接时优化（LTO）可以帮助减少分离编译的性能开销。

分离编译要求满足以下条件：

在一个编译单元中定义的非const设备变量，在其他编译单元中必须使用extern关键字引用。
所有const设备变量都必须使用extern关键字定义和引用。
所有 CUDA 源文件.cu都必须使用-dc或-rdc=true标志进行编译。

主机函数和设备函数默认具有外部链接性，不需要extern关键字。请注意，从 CUDA 13 开始，__global__函数和__managed__/__device__/__constant__变量默认具有内部链接性。

在以下示例中，definition.cu定义了一个变量和一个函数，而example.cu引用了它们。两个文件分别编译，然后链接成最终的二进制文件。

// ----- definition.cu -----
extern __device__ int device_variable = 5;
__device__        int device_function() { return 10; }

// ----- example.cu -----
extern __device__ int  device_variable;
__device__        int device_function();

__global__ void kernel(int* ptr) {
    device_variable = 0;
    *ptr            = device_function();
}

nvcc -dc definition.cu -o definition.o
nvcc -dc example.cu    -o example.o
nvcc definition.o example.o -o program

2.7.4. 常用编译器选项

本节介绍可以与nvcc配合使用的最相关的编译器选项，涵盖语言特性、优化、调试、性能分析和构建等方面。所有选项的完整说明可在nvcc 文档中找到。

2.7.4.1. 语言特性

nvcc支持从 C++03 到C++23 语言特性的 C++ 核心语言特性。-std标志可用于指定要使用的语言标准：

--std={c++03|c++11|c++14|c++17|c++20|c++23}

此外，nvcc还支持以下语言扩展：

-restrict：断言所有核函数指针参数都是restrict指针。
-extended-lambda：允许在 lambda 声明中使用__host__、__device__注解。
-expt-relaxed-constexpr：（实验性标志）允许主机代码调用__device__ constexpr函数，并允许设备代码调用__host__ constexpr函数。

有关这些特性的更多详细信息，请参阅扩展 lambda和constexpr部分。

2.7.4.2. 调试选项

nvcc支持以下用于生成调试信息的选项：

-g：为主机代码生成调试信息。gdb/lldb和类似工具依靠此类信息来调试主机代码。
-G：为设备代码生成调试信息。cuda-gdb依靠此类信息来调试设备代码。该标志还会定义__CUDACC_DEBUG__宏。
-lineinfo：为设备代码生成行号信息。此选项不影响执行性能，与compute-sanitizer工具配合使用时，有助于跟踪核函数的执行。

nvcc默认对 GPU 代码使用最高优化级别-O3。调试标志-G会阻止某些编译器优化，因此调试代码的性能预计会低于非调试代码。可以定义-DNDEBUG标志来禁用运行时断言，因为这些断言也可能降低执行速度。

2.7.4.3. 优化选项

nvcc提供许多用于优化性能的选项。本节旨在简要概述一些开发者可能会觉得有用的可用选项，并提供更多信息的链接。完整介绍可在nvcc 文档中找到。

-Xptxas将参数传递给 PTX 汇编器工具ptxas。nvcc文档提供了ptxas的实用参数列表。例如，-Xptxas=-maxrregcount=N指定每个线程可使用的最大寄存器数量。
-extra-device-vectorization：启用更激进的设备代码向量化。
--apply-controls=/path/to/file将高级控制文件（ACF）传递给 nvcc 和 ptxas。此文件会更改默认的编译行为，使其更有针对性地适应特定工作负载。使用高级控制文件可能导致编译失败或运行时执行不正确。使用风险自负。有关如何生成高级控制文件的更多信息，请参阅CompileIQ Github 页面。
有关对浮点行为提供细粒度控制的其他标志，请参阅浮点计算部分以及nvcc 文档。

以下标志可获取编译器的输出，这些输出对于更高级的代码优化可能很有用：

-res-usage：编译后打印资源使用报告。报告包括为每个核函数分配的寄存器数量以及共享内存、常量内存和局部内存的用量。
-opt-info=inline：打印有关内联函数的信息。
-Xptxas=-warn-lmem-usage：如果使用了局部内存，则发出警告。
-Xptxas=-warn-spills：如果寄存器溢出到局部内存，则发出警告。

2.7.4.4. 链接时优化（LTO）

由于跨文件优化的机会有限，分离编译可能导致性能低于全程序编译。链接时优化（LTO）通过在链接时对分别编译的文件进行跨文件优化来解决这个问题，代价是增加编译时间。LTO 可以在保留分离编译灵活性的同时，恢复全程序编译的大部分性能。

nvcc需要使用-dlto标志或lto_<SM version>链接时优化目标来启用 LTO：

nvcc -dc -dlto -arch=sm_100 definition.cu -o definition.o
nvcc -dc -dlto -arch=sm_100 example.cu    -o example.o
nvcc -dlto definition.o example.o -o program

nvcc -dc -arch=lto_100 definition.cu -o definition.o
nvcc -dc -arch=lto_100 example.cu    -o example.o
nvcc -dlto definition.o example.o -o program

2.7.4.5. 性能分析选项

可以使用Nsight Compute和Nsight Systems工具直接对 CUDA 应用程序进行性能分析，而无需在编译过程中添加额外的标志。不过，nvcc可以生成额外的信息，通过将源文件与生成的代码关联起来辅助性能分析：

-lineinfo：为设备代码生成行号信息；这允许在性能分析工具中查看源代码。性能分析工具要求原始源代码位于编译该代码时所在的同一位置。
-src-in-ptx：在 PTX 中保留原始源代码，避免上述-lineinfo的局限。需要-lineinfo。

2.7.4.6. Fatbin 压缩

nvcc默认会压缩存储在应用程序或库的二进制文件中的fatbin。可以使用以下选项控制 fatbin 压缩：

-no-compress：禁用 fatbin 压缩。
--compress-mode={default|size|speed|balance|none}：设置压缩模式。speed侧重于缩短解压缩时间，而size旨在减小 fatbin 大小。balance在速度和大小之间提供折中。默认模式为speed。none禁用压缩。

2.7.4.7. 编译器性能控制

nvcc提供了用于分析和加速编译过程本身的选项：

-t <N>：针对多个 GPU 架构并行编译单个编译单元时使用的 CPU 线程数。
-split-compile <N>：用于并行执行优化阶段的 CPU 线程数。
-split-compile-extended <N>：更激进的拆分编译形式。需要链接时优化。
-Ofc <N>：设备代码编译速度的级别。
-time <filename>：生成一个逗号分隔值（CSV）表，其中包含每个编译阶段所花费的时间。
-fdevice-time-trace：生成设备代码编译的时间跟踪记录。