0130W Notes

Article

CUDA笔记-NVCC

CUDA编译器nvcc相关内容

  • Notes
  • Computer
  • CUDA

nvcc是CUDA Toolkit提供的一个工具链, 用于将CUDA C/C++代码以及PTX代码编译为二进制文件, nvcc工具链中包含了多种工具, 包括compiler, linker以及PTX和Cubin的assembler. 顶层nvcc负责在对应的编译阶段唤醒相应的工具对源文件进行处理。

.cu文件为CUDA源文件，.cuh为CUDA头文件。 .cu和.cuh文件中可以包含host端以及device端上执行的代码。

NVCC Compilation Workflow

在初始化阶段， nvcc会先将代码文件中host端的代码与device端的代码区分开， 分别交给host端和device端对应的compiler处理。

对于仅包含host端的代码文件，也可以用例如clang等编译器编译出.o文件。

对于device端的代码部分，nvcc会先将源代码编译为PTX汇编代码，此时可以通过对compiler通过参数配置的方式指定需要输出的PTX代码的计算能力版本，例如compute_90。

生成的PTX代码部分会交给ptxas工具来生成指定SM版本的.cubin文件。生成的多个SM版本的.cubin文件和多个计算版本的PTX文件可以被打包到一个fatbin容器中， fatbin容器与host端可执行代码共同组成了最后的可执行文件。

通过传递-v参数给nvcc可以查看整个编译过程的工作流以及调用的工具。同时通过传递-keep(当前文件夹)和--keep-dir(指定对应文件夹)参数能够保留编译过程中的中间产物。

Example

代码部分:

// ----- 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;
}

编译过程Workflow示意图:

cuda-learn-nvcc-flow-multi-archs

NVCC Basic Usage

指定参数使用nvcc工具

nvcc example.cu \
-I <include_dir_path> \ # 哪里找header
-L <include_lib_path> \ # 哪里找libary
-l<library> \ #链接的library
-D<macro>=<value> \ # 定义宏，等价于在前面写#define macro value
-o <output_file>

NVCC PTX and Cubin Generation

默认不传参数的情况下，为了最大兼容性，nvcc会生成CUDA Toolkit能支持的最小compute版本号和sm版本号的PTX和Cubin文件。

如果需要指定生成目标GPU架构的PTX和cubin, 可以通过以下两种方式:

  1. -arch: 指定特定的一个GPU架构, 生成对应的PTX代码和Cubin文件

  2. -gencode: 指定多个GPU架构, 生成对应的PTX代码和Cubin文件

可以通过传递如下的参数给nvcc来查询支持的real GPU架构以及virtual GPU架构.

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

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

生成PTX代码和Cubin的指令类似下面:

nvcc example.cu -arch=compute_80 -gpu-code=sm_86

nvcc example.cu -gencode=arch=compute_80,code=sm_86,sm_89

更详细地说明, 可以参见nvcc documentation 会不会保留PTX, 主要是看-code/-gpu-code/code=后面有没有指定virtual GPU架构, 例如说compute_80等.

Host Code Compilation

如果Host端的代码依赖于CUDA Runtime API, 那么就应该链接CUDA Runtime library. CUDA Runtime提供静态库libcudart_static和动态库libcudart, 默认情况下nvcc会链接静态库, 可以传--cudart=shared要求nvcc链接动态库.

在编译Host端代码时, 可以指定对应的compiler, 有以下两种方式:

  1. -ccbin=<compiler>参数传递给nvcc, 例如-ccbin=clang++

  2. 设置环境变量NVCC_CCBIN=<compiler>, 例如export NVCC_CCBIN='gcc'

可以通过-XCompiler=<arguments>flag传递参数给对应的host compiler, 如-XCompiler=-O3

Separate Compilation of GPU Code

默认情况下, nvcc的编译方式为whole-program compilation, 即期望device端的符号和定义都在同一个编译单元(一个source file与其若干headers)中.

如果希望某个编译单元依赖于其他的编译单元, 则需要给nvcc传入-dc参数或-rdc=true参数用于开启链接时能够跨编译单元进行链接. 这种编译方式称为separate compilation.

separate compilation会带来一些运行时影响, 这是因为在whole-program compilation的情况下, nvcc能够一次看到全部的device代码. 如果对性能比较敏感, 可以对nvcc传入-dto参数或指定-arch=lto_<SM version>来开启Link-Time Optimization(LTO), 通过增加额外的编译时间, 做跨文件的链接时优化, 使得separate compilation的性能接近于whole-program compilation的编译方式.

在CUDA 13之后, __global__函数和__managed__/__device__/__constant__变量在a.cu中定义, 并在b.cu中引用时, 必须要加上extern关键字, 这是因为CUDA 13之后, 这些都是默认内部链接, 对外不可见的. __device__函数普通的c++符号继续沿用c++习惯无需extern关键字.

Common Compiler Options

Languages Features

nvcc支持从c++03到c++23的标准, 可以通过--std参数指定, --std={c++03|c++11|c++14|c++17|c++20|c++23}

支持以下的常见flag:

  1. -restrict: 标明所有的kernel的指针入参都是restrict的

  2. -extended-lambda: 支持在使用lambda表达式时加入__host__和__device__记号, 见extended lambda

  3. -expt-relaxed-constexpr(Experimental flag): 使得host端的函数能够调用device端的constexpr函数(允许编译时计算), device端的函数能够调用host端的constexpr函数., 见constexpr.

CUDA官方不推荐在__device__端使用c++标准库std::, 建议使用cuda::std::中的函数.

Debugging Options

  1. -g: 生成host端debug信息, 后续可以给gdb/lldb之类的工具分析

  2. -G: 生成device端debug信息, 同时定义__CUDACC_DEBUG__宏, 用于给cuda-gdb工具分析, 由于默认情况下对device端代码进行编译的时候用的是-O3级别优化, 开启-G参数可能导致运行变慢.

  3. -lineinfo: 为device端代码生成行号信息, 不影响运行时性能, 可以给compute-sanitizer工具来追踪kernel执行情况

  4. -DNDEBUG: 关闭运行时断言, 可以加速运行, 一般可以在release版本中加.

Optimization Options

  1. -Xptxas=<args>: 将参数传给PTX assembler tool ptxas, 具体见ptxas arguments

  2. -extra-device-vectorization: 允许更激进的device端代码向量化

3.--apply-controls=/path/to/file: 将advanced controls file(ACF)传递给nvcc以及ptxas, 改变编译时的行为, 例如说kernel应该怎么内联, 怎么重排等. 这个可能会导致编译失败, 一般推荐用CompileIQ生成ACF.

  1. 一些控制浮点数行为的flag, 见Floating-Point Computation以及nvcc documentation

前面separate compilation部分已经提过.

Profilling Options

使用Nsight Compute和Nsight Systems工具可以对CUDA程序进行性能分析, 添加如下的参数可以将原文件和生成的文件进行关联

  1. -lineinfo: 为device端代码生成行号信息, 从而可以在profiling工具中直接看到源代码

  2. -src-in-ptx: 在生成的PTX代码中包含源代码, 需要开启-lineinfo.

Fatbin Compression

默认情况下, nvcc会对fatbins做压缩, 可以通过如下的参数来对压缩做进一步的管理:

  1. -no-compression: 关闭对fatbin的压缩.

  2. --compression-mode={default|size|speed|balance|none}: default为默认行为, size追求更小的fatbin体积, speed追求更快的解压缩速度, balance平衡size和speed的优化, none关闭对fatbin的压缩.

Compiler Performance Controls

nvcc提供了一些参数用于分析和加速编译过程.

  1. -t <N>: 指定用于编译生成多GPU架构的单个编译单元所用的线程数

  2. -time <filename>: 生成各编译阶段的编译耗时, 并存放在CSV文件中

  3. -fdevice-time-trace: 对device端代码的编译过程时间进行trace.

  4. -Ofc <N>: 允许max/mid/min/0四档, 0为关闭fast compile, 本质是拿生成代码的运行性能换编译速度, 不建议开.

  5. -split-compile <N>: 允许ptxas将优化阶段拆分为多个模块, 用N个CPU线程并行处理.

  6. -split-compile-extended <N>: 类似于-split-compile, 不过将处理推迟到了link阶段, 此时拿到的是全局device代码, 可以做更激进的编译速度优化, 需要开启-dlto.

整理自NVCC: The NVIDIA CUDA Compiler