NVIDIA CUDA 编译器驱动 NVCC
nvcc的文档，CUDA编译器驱动程序。
1. 简介
1.1. 概述
1.1.1. CUDA编程模型
CUDA工具包面向一类应用程序，其控制部分作为进程运行在通用计算设备上，并使用一个或多个NVIDIA GPU作为协处理器来加速单程序多数据(SPMD)并行任务。这类任务是自包含的，即它们可以完全由一批GPU线程执行和完成，无需主机进程干预，从而从并行图形硬件中获得最佳效益。
GPU代码以函数集合的形式实现，使用的语言本质上是C++，但带有一些注解以区分主机代码，以及用于区分GPU上不同类型数据内存的标注。这些函数可以包含参数，调用语法与常规C函数调用非常相似，但略有扩展以便能够指定必须执行被调用函数的GPU线程矩阵。在主机进程的生命周期内，可以分派许多并行GPU任务。
如需了解更多关于CUDA编程模型的信息，请参阅CUDA C++ Programming Guide。
1.1.2. CUDA 源代码
CUDA应用程序的源文件由传统的C++主机代码和GPU设备函数混合组成。CUDA编译流程将设备函数与主机代码分离，使用专有的NVIDIA编译器和汇编器编译设备函数，同时使用可用的C++主机编译器编译主机代码，随后将编译后的GPU函数作为fatbinary镜像嵌入主机目标文件中。在链接阶段，会添加特定的CUDA运行时库，以支持远程SPMD过程调用，并提供显式的GPU操作，例如GPU内存缓冲区的分配和主机-GPU数据传输。
1.1.3. NVCC的目的
编译轨迹涉及对每个CUDA源文件进行多次分割、编译、预处理和合并步骤。nvcc作为CUDA编译器驱动程序，其目的是向开发者隐藏CUDA编译的复杂细节。它接受一系列常规编译器选项，例如用于定义宏和包含/库路径，以及控制编译过程。所有非CUDA编译步骤都会被转发到nvcc支持的C++主机编译器，nvcc会将其选项转换为适当的主机编译器命令行选项。
1.2. 支持的宿主编译器
在以下情况下，nvcc需要一个通用的C++主机编译器：
在非CUDA阶段（运行阶段除外），因为这些阶段将由
nvcc转发给此编译器。在CUDA阶段，用于多个预处理阶段和主机代码编译（另请参阅CUDA编译轨迹）。
nvcc 假设主机编译器已按照编译器供应商设计的标准方法安装。如果主机编译器安装方式非标准，用户必须确保环境设置正确，并使用相关的 nvcc 编译选项。
以下文档提供了关于支持的主机编译器的详细信息：
在所有平台上，除非通过适当选项另行指定（参见file-and-path-specifications），否则将使用当前执行搜索路径中找到的默认主机编译器可执行文件（Linux上为gcc和g++，Windows上为cl.exe）。
注意，nvcc不支持编译超过主机系统最大路径长度限制的文件路径。如需支持编译长文件路径，请参阅您系统的相关文档。
2. 编译阶段
2.1. NVCC 标识宏
nvcc 预定义了以下宏：
__NVCC__在编译C/C++/CUDA源文件时定义。
__CUDACC__在编译CUDA源文件时定义。
__CUDACC_RDC__在可重定位设备代码模式下编译CUDA源文件时定义（参见NVCC独立编译选项）。
__CUDACC_EWP__在可扩展全程序模式下编译CUDA源文件时定义（参见指定编译器/链接器行为的选项）。
__CUDACC_DEBUG__在设备调试模式下编译CUDA源文件时定义（参见指定编译器/链接器行为的选项）。
__CUDACC_RELAXED_CONSTEXPR__当在命令行中指定
--expt-relaxed-constexpr标志时定义。更多详情请参阅CUDA C++编程指南。__CUDACC_EXTENDED_LAMBDA__当在命令行中指定
--expt-extended-lambda或--extended-lambda标志时定义。更多详情请参阅CUDA C++ Programming Guide。__CUDACC_VER_MAJOR__由
nvcc的主版本号定义。__CUDACC_VER_MINOR__由
nvcc的次版本号定义。__CUDACC_VER_BUILD__使用
nvcc的构建版本号定义。__NVCC_DIAG_PRAGMA_SUPPORT__当CUDA前端编译器支持使用
nv_diag_suppress、nv_diag_error、nv_diag_warning、nv_diag_default、nv_diag_once和nv_diagnostic编译指示进行诊断控制时定义。__CUDACC_DEVICE_ATOMIC_BUILTINS__当CUDA前端编译器支持设备原子编译器内置函数时定义。更多详情请参阅CUDA C++编程指南。
2.2. NVCC编译阶段
A compilation phase is a logical translation step that can be selected by command line options to nvcc. A single compilation phase can still be broken up by nvcc into smaller steps, but these smaller steps are just implementations of the phase: they depend on seemingly arbitrary capabilities of the internal tools that nvcc uses, and all of these internals may change with a new release of the CUDA Toolkit. Hence, only compilation phases are stable across releases, and although nvcc provides options to display the compilation steps that it executes, these are for debugging purposes only and must not be copied and used in build scripts.
nvcc 阶段的选择由命令行选项和输入文件后缀共同决定，这些阶段的执行可能会被其他命令行选项修改。在阶段选择中，输入文件后缀定义了阶段的输入，而命令行选项则定义了该阶段所需的输出。
以下段落列出了可识别的文件扩展名及支持的编译阶段。关于nvcc命令行选项的完整说明，请参阅NVCC Command Options。
2.3. 支持的输入文件后缀
下表定义了nvcc如何解析其输入文件：
输入文件后缀 |
描述 |
|---|---|
|
CUDA源文件，包含主机代码和设备函数 |
|
C源文件 |
|
C++源文件 |
|
PTX中间汇编文件（参见图1） |
|
针对单一GPU架构的CUDA设备代码二进制文件(CUBIN)（参见图1） |
|
CUDA fat二进制文件，可能包含多个PTX和CUBIN文件（参见：图1 |
|
目标文件 |
|
库文件 |
|
资源文件 |
|
共享对象文件 |
请注意，nvcc不会区分对象文件、库文件或资源文件。在执行链接阶段时，它只是将这些类型的文件传递给链接器。
2.4. 支持的阶段
下表列出了支持的编译阶段，以及启用每个阶段执行的nvcc选项。同时还列出了每个阶段生成的输出文件默认名称，当未使用--output-file选项显式指定输出文件名时，这些默认名称将生效：
阶段 |
|
默认输出文件名 |
|
|---|---|---|---|
长名称 |
短名称 |
||
CUDA编译为C/C++源文件 |
|
|
|
C/C++预处理 |
|
|
<标准输出结果> |
C/C++ 编译为目标文件 |
|
|
源文件名后缀在Linux上替换为 |
从CUDA源文件生成Cubin |
|
|
源文件名后缀替换为 |
从PTX中间文件生成Cubin。 |
|
|
源文件名后缀替换为 |
从CUDA源文件生成PTX |
|
|
源文件名后缀替换为 |
从源代码、PTX或cubin文件生成Fatbinary |
|
|
源文件名后缀替换为 |
链接可重定位设备代码。 |
|
|
|
从链接的可重定位设备代码生成Cubin。 |
|
|
|
从链接的可重定位设备代码生成Fatbinary |
|
|
|
链接可执行文件 |
<无阶段选项> |
|
|
构建一个对象文件归档或库 |
|
|
|
|
|
|
<标准输出结果> |
|
|
|
<标准输出结果> |
将CUDA源代码编译为OptiX IR输出。 |
|
|
源文件名后缀替换为 |
将CUDA源代码编译为LTO IR输出。 |
|
|
源文件名后缀替换为 |
运行可执行文件 |
|
|
注意：
列表中的最后一个阶段更多是为了方便。它允许运行已编译和链接的可执行文件，而无需显式设置CUDA动态库的库路径。
除非指定了阶段选项，否则
nvcc将编译并链接其所有输入文件。
3. CUDA编译流程
CUDA的编译过程如下：输入程序首先经过设备编译预处理，被编译为CUDA二进制文件(cubin)和/或PTX中间代码，这些内容会被放入一个fatbinary中。随后输入程序会再次进行主机编译预处理，通过代码合成将fatbinary嵌入其中，并将CUDA特定的C++扩展转换为标准C++结构。最后，C++主机编译器会将嵌入了fatbinary的合成主机代码编译成主机对象文件。具体步骤如图1所示。
每当主机程序启动设备代码时，CUDA运行时系统会检查嵌入的fatbinary，以获取适用于当前GPU的fatbinary镜像。
CUDA程序默认采用全程序编译模式，即设备代码无法引用单独文件中的实体。在全程序编译模式下，设备链接步骤不会产生任何效果。有关单独编译和全程序编译的更多信息，请参阅Using Separate Compilation in CUDA。
CUDA编译轨迹
4. NVCC 命令选项
4.1. 命令选项类型与表示法
每个nvcc选项都有一个长名称和一个短名称，两者可以互换使用。这两种变体的区别在于选项名称前需要添加的连字符数量：长名称前需要两个连字符，而短名称前只需一个连字符。例如，-I是--include-path的短名称。长选项适用于构建脚本，此时选项的描述性价值比长度更重要；而短选项则更适合交互式使用。
nvcc 识别三种类型的命令行选项：布尔选项、单值选项和列表选项。
布尔选项不需要参数；它们要么在命令行中指定，要么不指定。单值选项最多只能指定一次，但列表选项可以重复。这些选项类型的示例分别是：--verbose（切换到详细模式）、--output-file（指定输出文件）和--include-path（指定包含路径）。
单值选项和列表选项必须带有参数，这些参数必须紧跟在选项名称之后，通过一个或多个空格或等号字符分隔。当使用单字符短名称如-I、-l和-L时，选项的值也可以直接跟在选项后面，无需通过空格或等号分隔。列表选项的各个值可以通过逗号在单个选项实例中分隔，也可以重复该选项，或者这两种情况的任意组合。
因此，对于上述两个可能取值的样本选项，以下表示方式是合法的：
-o file
-o=file
-Idir1,dir2 -I=dir3 -I dir4,dir5
除非另有说明，本文档中均使用长选项名称。但短名称可以替代长名称，效果相同。
4.2. 命令选项说明
本节展示nvcc选项的表格。表格中的选项类型可通过以下方式识别：布尔选项在第一列没有指定参数，而其他两种类型则有。列表选项可通过参数末尾的重复指示符,...来识别。
长选项在选项表格的第一列中描述，短选项位于第二列。
4.2.1. 文件和路径规范
4.2.1.1. --output-file file (-o)
指定输出文件的名称和位置。
4.2.1.2. --objdir-as-tempdir (-objtemp)
在与目标文件相同的目录中创建所有中间文件。这些中间文件在编译完成后会被删除。此选项仅在同时使用-c、-dc或-dw时生效。使用此选项可确保嵌入目标文件的中间文件名在多次编译同一文件时不会改变。但如果输入来自stdin则无法保证此效果。若同一文件使用两种不同选项编译（例如'nvcc -c t.cu'和'nvcc -c -ptx t.cu'），则应在不同目录中进行编译。在同一目录中编译可能导致编译失败或产生错误结果。
4.2.1.3. --pre-include file,... (-include)
指定在预处理阶段必须预先包含的头文件。
4.2.1.4. --library library,... (-l)
指定在链接阶段使用的库，无需包含库文件扩展名。
库文件会在通过选项--library-path指定的库搜索路径中进行查找（参见Libraries）。
4.2.1.5. --define-macro def,... (-D)
定义在预处理过程中使用的宏。
def 可以是 name 或 name=definition。
name - 将name预定义为宏。
name=definition - The contents of definition are tokenized and preprocessed as if they appear during translation phase three in a
#definedirective. The definition will be truncated by embedded new line characters.
4.2.1.6. --undefine-macro def,... (-U)
在预处理或编译期间取消定义现有的宏。
4.2.1.7. --include-path path,... (-I)
指定包含搜索路径。
4.2.1.8. --system-include path,... (-isystem)
指定系统包含的搜索路径。
4.2.1.9. --library-path path,... (-L)
指定库搜索路径（参见Libraries）。
4.2.1.10. --output-directory directory (-odir)
指定输出文件的目录。
此选项旨在让依赖生成步骤（参见--generate-dependencies）生成一个规则，该规则在正确的目录中定义目标对象文件。
4.2.1.11. --dependency-output file (-MF)
指定依赖项输出文件。
此选项指定依赖生成步骤的输出文件（参见--generate-dependencies）。如果设置了依赖输出文件，则必须指定选项--generate-dependencies或--generate-nonystem-dependencies。
4.2.1.12. --generate-dependency-targets (-MP)
为每个依赖项添加一个空目标。
此选项会在依赖生成步骤中添加伪目标（参见--generate-dependencies），旨在避免当旧依赖被删除时出现makefile错误。输入文件不会作为伪目标输出。
4.2.1.13. --compiler-bindir directory (-ccbin)
指定默认主机编译器可执行文件所在的目录。
可以指定主机编译器可执行文件的名称，以确保选择正确的主机编译器。此外，如果在Windows的Cygwin shell或MinGW shell中执行nvcc，可能需要指定驱动前缀选项（--input-drive-prefix、--dependency-drive-prefix或--drive-prefix）。
4.2.1.14. --allow-unsupported-compiler (-allow-unsupported-compiler)
禁用对受支持主机编译器版本的nvcc检查。
使用不受支持的主机编译器可能导致编译失败或运行时执行错误。风险自负。此选项在MacOS上无效。
4.2.1.15. --archiver-binary executable (-arbin)
指定用于创建静态库的归档工具路径，使用--lib参数。
4.2.1.17. --cudadevrt {none|static} (-cudadevrt)
指定要使用的CUDA设备运行时库类型：不使用CUDA设备运行时库，或静态CUDA设备运行时库。
允许的值
nonestatic
默认
默认使用静态CUDA设备运行时库。
4.2.1.18. --libdevice-directory directory (-ldir)
指定包含libdevice库文件的目录。
Libdevice库文件位于CUDA工具包中的nvvm/libdevice目录下。
4.2.1.19. --target-directory string (-target-dir)
指定目标目录中的子文件夹名称，其中包含默认的包含路径和库路径。
4.2.2. 编译阶段指定选项
此类别选项指定输入文件必须编译到哪个阶段。
4.2.2.1. --link (-link)
指定默认行为：编译并链接所有输入文件。
默认输出文件名
a.exe 在Windows系统上或 a.out 在其他平台上被用作默认的输出文件名。
4.2.2.2. --lib (-lib)
如有必要，将所有输入文件编译为目标文件，并将结果添加到指定的库输出文件中。
默认输出文件名
a.lib 在Windows平台上或 a.a 在其他平台上被用作默认输出文件名。
4.2.2.3. --device-link (-dlink)
将包含可重定位设备代码的目标文件与.ptx、.cubin和.fatbin文件链接成一个包含可执行设备代码的目标文件，该文件可以传递给主机链接器。
默认输出文件名
a_dlink.obj 在Windows平台上或 a_dlink.o 在其他平台上被用作默认输出文件名。当此选项与 --fatbin 一起使用时，a_dlink.fatbin 将作为默认输出文件名。当此选项与 --cubin 一起使用时，a_dlink.cubin 将作为默认输出文件名。
4.2.2.4. --device-c (-dc)
将每个.c、.cc、.cpp、.cxx和.cu输入文件编译成包含可重定位设备代码的目标文件。
它等同于 --relocatable-device-code=true --compile。
默认输出文件名
源文件扩展名在Windows上被替换为.obj，在其他平台上替换为.o以创建默认输出文件名。例如，x.cu的默认输出文件名在Windows上是x.obj，在其他平台上是x.o。
4.2.2.5. --device-w (-dw)
将每个.c、.cc、.cpp、.cxx和.cu输入文件编译成包含可执行设备代码的目标文件。
它等同于 --relocatable-device-code=false --compile。
默认输出文件名
源文件扩展名在Windows上被替换为.obj，在其他平台上替换为.o以创建默认输出文件名。例如，x.cu的默认输出文件名在Windows上是x.obj，在其他平台上是x.o。
4.2.2.6. --cuda (-cuda)
将每个.cu输入文件编译为.cu.cpp.ii文件。
默认输出文件名
.cu.cpp.ii会附加在源文件名的基本名称后，作为默认的输出文件名。例如，x.cu的默认输出文件名是x.cu.cpp.ii。
4.2.2.7. --compile (-c)
将每个.c、.cc、.cpp、.cxx和.cu输入文件编译成目标文件。
默认输出文件名
源文件扩展名在Windows上被替换为.obj，在其他平台上替换为.o以创建默认输出文件名。例如，x.cu的默认输出文件名在Windows上是x.obj，在其他平台上是x.o。
4.2.2.8. --fatbin (-fatbin)
将所有.cu、.ptx和.cubin输入文件编译为仅限设备的.fatbin文件。
nvcc 使用此选项时会丢弃每个.cu输入文件的主机代码。
默认输出文件名
源文件名的扩展名会被替换为.fatbin以生成默认输出文件名。例如，x.cu的默认输出文件名是x.fatbin。
4.2.2.9. --cubin (-cubin)
将所有.cu和.ptx输入文件编译为仅限设备的.cubin文件。
nvcc 使用此选项时会丢弃每个.cu输入文件的主机代码。
默认输出文件名
源文件的扩展名会被替换为.cubin以生成默认的输出文件名。例如，x.cu的默认输出文件名是x.cubin。
4.2.2.10. --ptx (-ptx)
将所有.cu输入文件编译为仅设备的.ptx文件。
nvcc 使用此选项时会丢弃每个.cu输入文件的主机代码。
默认输出文件名
源文件的扩展名会被替换为.ptx以生成默认的输出文件名。例如，x.cu的默认输出文件名是x.ptx。
4.2.2.11. --preprocess (-E)
预处理所有.c、.cc、.cpp、.cxx和.cu输入文件。
默认输出文件名
默认情况下，输出会生成在stdout中。
4.2.2.12. --generate-dependencies (-M)
生成一个可包含在Makefile中的依赖文件，适用于.c、.cc、.cpp、.cxx和.cu输入文件。
nvcc 使用固定前缀来识别预处理文件中的依赖项（Linux上是‘#line 1’，Windows上是‘# 1’）。以该前缀开头的源位置指令中提到的文件将被包含在依赖列表中。
默认输出文件名
默认情况下，输出会显示在stdout中。
4.2.2.13. --generate-nonsystem-dependencies (-MM)
与--generate-dependencies相同，但会跳过在系统目录中找到的头文件（仅限Linux）。
默认输出文件名
默认情况下，输出会显示在stdout中。
4.2.2.14. --generate-dependencies-with-compile (-MD)
生成依赖文件并编译输入文件。该依赖文件可被包含在Makefile中，适用于.c、.cc、.cpp、.cxx和.cu格式的输入文件。
此选项不能与 -E 同时指定。依赖项文件名的计算方式如下：
如果指定了
-MF，则使用指定文件作为依赖文件名。如果指定了
-o选项，依赖文件名将通过将指定文件名的后缀替换为'.d'来计算得出。否则，依赖文件名将通过把输入文件名的后缀替换为'.d'来计算得出。
如果依赖文件名是基于-MF或-o计算得出的，则不支持多个输入文件。
4.2.2.15. --generate-nonsystem-dependencies-with-compile (-MMD)
与--generate-dependencies-with-compile相同，但会跳过在系统目录中找到的头文件（仅限Linux系统）。
4.2.2.16. --optix-ir (-optix-ir)
将CUDA源代码编译为OptiX IR (.optixir)输出。OptiX IR仅适用于通过适当API由OptiX使用。此功能不支持链接时优化(-dlto)、lto_NN架构目标或-gencode。
默认输出文件名
源文件名的扩展名被替换为.optixir以创建默认的输出文件名。例如，x.cu的默认输出文件名是x.optixir。
4.2.2.17. --ltoir (-ltoir)
将CUDA源代码编译为LTO IR (.ltoir)输出。此功能仅支持链接时优化(-dlto)或lto_NN架构目标。
默认输出文件名
源文件名的扩展名将被替换为.ltoir以创建默认的输出文件名。例如，x.cu的默认输出文件名是x.ltoir。
4.2.2.18. --run (-run)
将所有输入文件编译并链接成一个可执行文件，然后执行它。
当输入是一个单独的可执行文件时，它会被直接执行而无需任何编译或链接步骤。这一步骤旨在为那些不想费心设置必要环境变量的开发者提供便利；这些环境变量将由nvcc临时设置。
4.2.3. 编译器/链接器行为指定选项
4.2.3.1. --profile (-pg)
为gprof工具生成可插桩的代码/可执行文件。
4.2.3.2. --debug (-g)
为主机代码生成调试信息。
4.2.3.3. --device-debug (-G)
为设备代码生成调试信息。
如果未指定--dopt，则此选项会关闭设备代码的所有优化。该选项不适用于性能分析；进行性能分析时请改用--generate-line-info。
4.2.3.4. --extensible-whole-program (-ewp)
生成可扩展的完整程序设备代码，允许某些调用在链接libcudadevrt之前保持未解析状态。
4.2.3.5. --no-compress (-no-compress)
不要在fatbinary中压缩设备代码。
4.2.3.6. --compress-mode {default|size|speed|balance|none} (-compress-mode)
选择fatbinary中的设备代码压缩行为。
允许的值
default
使用默认压缩模式，如同未指定此选项。此模式的行为可能随版本变化而改变。目前等同于
speed模式。
size
采用一种更侧重于减小二进制文件大小的压缩模式，但会牺牲压缩和解压缩时间。
speed
采用一种更注重减少解压时间的压缩模式，代价是最终二进制文件大小的缩减幅度较小。
balance
采用一种平衡二进制大小与压缩/解压时间的压缩模式。
none
不执行压缩。等同于
--no-compress。
默认值
default 被用作默认模式。
4.2.3.7. --generate-line-info (-lineinfo)
为设备代码生成行号信息。
4.2.3.8. --optimization-info kind,... (-opt-info)
为指定类型的优化提供优化报告。
支持以下标签：
inline
输出与函数内联相关的备注信息。编译器可能会多次调用内联过程，在前一轮未内联的函数可能在后续轮次中被内联。
4.2.3.9. --optimize level (-O)
为主机代码指定优化级别。
4.2.3.10. --dopt kind (-dopt)
启用设备代码优化。当与-G一起指定时，可为优化后的设备代码生成有限的调试信息（目前仅支持行号信息）。当未指定-G时，-dopt=on将默认启用。
允许的值
on: 启用设备代码优化。
4.2.3.11. --dlink-time-opt (-dlto)
对设备代码执行链接时优化。选项'-lto'也是'-dlto'的别名。链接时优化必须在编译和链接时同时指定：在编译时存储高级中间代码，然后在链接时将这些中间代码链接在一起并进行优化。如果在链接时未找到该中间代码，则不会发生任何操作。中间代码在编译时也会通过--gpu-code='lto_NN'目标进行存储。选项-dlto -arch=sm_NN将添加一个lto_NN目标；如果只想添加lto_NN目标而不添加-arch=sm_NN通常生成的compute_NN目标，请使用-arch=lto_NN。
4.2.3.12. --gen-opt-lto (-gen-opt-lto)
在生成LTO IR之前运行优化器过程。
4.2.3.13. --split-compile number (-split-compile)
并行执行编译器优化。
Split compilation attempts to reduce compile time by enabling the compiler to run certain optimization passes concurrently. It does this by splitting the device code into smaller translation units, each containing one or more device functions, and running optimization passes on each unit concurrently across multiple threads. It will then link back the split units prior to code generation.
The option accepts a numerical value that specifies the maximum number of threads the compiler can use. One can also allow the compiler to use the maximum threads available on the system by setting --split-compile=0. Setting --split-compile=1 will cause this option to be ignored.
This option can work in conjunction with device Link Time Optimization (-dlto) as well as --threads.
4.2.3.14. --split-compile-extended number (-split-compile-extended)
一种更为激进的-split-compile形式。仅在LTO模式下可用。
扩展分割编译通过将并发编译延伸到后端，试图进一步减少编译时间。这种激进的分割编译形式可能会影响编译后二进制文件的性能。
该选项接受一个数值，用于指定编译器可以使用的最大线程数。也可以通过设置--split-compile-extended=0允许编译器使用系统上可用的最大线程数。设置--split-compile-extended=1将导致此选项被忽略。
此选项仅适用于设备链接时优化(-dlto)，并且可以与--threads一起使用。
4.2.3.15. --ftemplate-backtrace-limit limit (-ftemplate-backtrace-limit)
设置单个警告或错误的模板实例化注释的最大数量限制。
允许值为0，表示不强制执行任何限制。如果主机编译器提供等效标志，此值也会传递给主机编译器。
4.2.3.16. --ftemplate-depth limit (-ftemplate-depth)
设置模板类的最大实例化深度以进行限制。
如果主机编译器提供了等效的标志，该值也会传递给主机编译器。
4.2.3.17. --no-exceptions (-noeh)
禁用主机代码的异常处理。
通过在主编译器调用期间传递“-EHs-c-”（针对cl.exe）和“–fno-exceptions”（针对其他主编译器），禁用主机代码的异常处理。这些标志会在使用“-Xcompiler”直接传递给主编译器的任何标志之前，添加到主编译器调用中。
默认（在Windows上）
在Windows系统上，
nvcc默认会将/EHsc参数传递给主机编译器。
示例（Windows系统）
nvcc --no-exceptions -Xcompiler /EHa x.cu
4.2.3.19. --x {c|c++|cu} (-x)
明确指定输入文件的语言，而不是让编译器根据文件后缀名选择默认语言。
允许的值
cc++cu
默认
源代码的语言是根据文件名的后缀来确定的。
4.2.3.20. --std {c++03|c++11|c++14|c++17|c++20} (-std)
选择特定的C++方言。
允许的值
c++03c++11c++14c++17c++20
默认
默认的C++方言取决于主机编译器。nvcc会匹配主机编译器使用的默认C++方言。
4.2.3.21. --no-host-device-initializer-list (-nohdinitlist)
不要将std::initializer_list的成员函数视为隐式的__host__ __device__函数。
4.2.3.22. --expt-relaxed-constexpr (-expt-relaxed-constexpr)
实验性标志: 允许主机代码调用``__device__ constexpr``函数，设备代码调用``__host__ constexpr``函数。
请注意，此标志的行为在未来编译器版本中可能会发生变化。
4.2.3.23. --extended-lambda (-extended-lambda)
允许在lambda声明中使用__host__、__device__注解。
4.2.3.24. --expt-extended-lambda (-expt-extended-lambda)
--extended-lambda的别名。
4.2.3.25. --machine {64} (-m)
指定64位架构。
允许的值
64
默认
此选项根据执行nvcc的主机平台设置。
4.2.3.26. --m64 (-m64)
--machine=64的别名
4.2.3.27. --host-linker-script {use-lcs|gen-lcs} (-hls)
使用主机链接器脚本（仅限GNU/Linux系统）来支持某些CUDA特定需求，同时构建可执行文件或共享库。
允许的值
use-lcs
准备一个主机链接器脚本，并启用主机链接器以支持体积较大的可重定位设备对象文件。在某些情况下，这些文件原本会导致主机链接器因重定位截断错误而失败。
gen-lcs
生成一个主机链接器脚本，可在主机链接器在nvcc之外单独调用时手动传递给主机链接器。此选项可与
-shared或-r选项结合使用，分别生成用于创建主机共享库或主机可重定位链接的链接器脚本。使用此选项生成的文件必须作为主机链接器的最后一个输入文件提供。
默认情况下，输出会生成到stdout。使用选项
-ofilename可以指定输出文件名。
可能已有链接脚本正在使用，并通过主机链接器选项--script（或-T）传递给主机链接器，此时生成的链接脚本必须对现有链接脚本进行扩展。在这种情况下，必须使用选项-aug-hls来生成仅包含扩展部分的链接脚本。否则，主机链接器的行为将无法确定。
一个主机链接器选项，例如带有非默认参数的-z，可能会在内部修改默认链接器脚本，与该选项不兼容，任何此类使用的行为都是未定义的。
默认值
use-lcs 被用作默认类型。
4.2.3.28. --augment-host-linker-script (-aug-hls)
启用生成主机链接器脚本的功能，该脚本可增强现有的主机链接器脚本（仅限GNU/Linux系统）。详情请参阅选项--host-linker-script。
4.2.3.29. --host-relocatable-link (-r)
当与-hls=gen-lcs组合使用时，控制-hls=gen-lcs的行为，将其设置为生成可用于主机可重定位链接的主机链接器脚本(ld -r链接)。更多信息请参阅选项-hls=gen-lcs。
此选项目前仅在与-hls=gen-lcs一起使用时有效；在所有其他情况下，该选项目前会被忽略。
4.2.4. 传递特定阶段选项的设置
这些标志允许直接将特定选项传递给nvcc封装的内部编译工具，而无需让nvcc了解这些工具的过多细节。
4.2.4.1. --compiler-options options,... (-Xcompiler)
直接向编译器/预处理器指定选项。
4.2.4.2. --linker-options options,... (-Xlinker)
直接向主机链接器指定选项。
4.2.4.3. --archive-options options,... (-Xarchive)
直接向库管理器指定选项。
4.2.4.4. --ptxas-options options,... (-Xptxas)
直接向ptxas（PTX优化汇编器）指定选项。
4.2.4.5. --nvlink-options options,... (-Xnvlink)
直接指定选项给nvlink，即设备链接器。
4.2.5. 指导编译器驱动的选项
4.2.5.1. --static-global-template-stub {true|false} (-static-global-template-stub)
在全程序编译模式下（-rdc=false），强制为__global__函数模板生成的宿主端存根函数使用static静态链接。
A __global__ function represents the entry point for GPU code execution, and is typically referenced from host code. In whole program compilation mode (nvcc default), the device code in each translation unit forms a self-contained device program. In the code sent to the host compiler, the CUDA frontend compiler will replace the contents of the body of the original __global__ function or function template with calls to the CUDA runtime to launch the kernel (these are referred to as ‘stub’ functions below).
When this flag is false, the template stub function will have weak linkage. This causes a problem if two different translation units a.cu and b.cu have the same instatiation for a __global__ template G.
例如：
//common.h
template <typename T>
__global__ void G() { qqq = 4; }
//a.cu
static __device__ int qqq;
#include "common.h"
int main() { G<int><<<1,1>>>(); }
//b.cu
static __device__ int qqq;
#include "common.h"
int main() { G<int><<<1,1>>>(); }
当a.cu和b.cu在nvcc全程序模式下编译时，为a.cu和b.cu生成的设备程序是独立的，但主机链接器会遇到G存根实例化的多个弱定义，并在链接的主机程序中仅选择其中一个。因此，从a.cu或b.cu启动G将错误地启动对应于a.cu或b.cu之一的设备程序；而正确的预期行为应该是：从a.cu启动的G运行为a.cu生成的设备程序，从b.cu启动的G运行为b.cu生成的设备程序。
当该标志为true时，CUDA前端编译器会在生成的主机代码中将所有存根函数设为static。这解决了上述问题，因为a.cu和b.cu中的G现在指向主机目标代码中的不同符号，主机链接器将不会合并这些符号。
注意事项
除非程序以全程序编译模式(
-rdc=false)进行编译，否则此选项将被忽略。启用此标志可能会在某些极端情况下破坏现有代码（仅在完整程序编译模式下）：
如果一个
__global__函数模板被声明为友元，且该友元声明是该实体的首次声明。If a
__global__function template is referenced, but not defined in the current translation unit.
默认
false。请注意，在未来的CUDA工具包主要版本中，默认值将更改为true。
4.2.5.3. --forward-unknown-to-host-compiler (-forward-unknown-to-host-compiler)
将未知选项传递给主机编译器。所谓"未知选项"是指以-开头后跟其他字符的命令行参数，且不被识别为nvcc标志或已知nvcc标志的参数。
如果未知选项后跟一个单独的命令行参数，除非该参数以-字符开头，否则该参数将不会被转发。
例如：
nvcc -forward-unknown-to-host-compiler -foo=bar a.cu会将-foo=bar转发给主机编译器。nvcc -forward-unknown-to-host-compiler -foo bar a.cu将会对bar参数报错。nvcc -forward-unknown-to-host-compiler -foo -bar a.cu会将-foo和-bar转发给主机编译器。
注意：在Windows系统上，另请参阅选项-forward-slash-prefix-opts以转发以'/'开头的选项。
4.2.5.4. --forward-unknown-to-host-linker (-forward-unknown-to-host-linker)
将未知选项转发给主机链接器。所谓"未知选项"是指以-开头后跟其他字符的命令行参数，且不被识别为nvcc标志或已知nvcc标志的参数。
如果未知选项后跟一个单独的命令行参数，除非该参数以-字符开头，否则该参数不会被转发。
例如：
nvcc -forward-unknown-to-host-linker -foo=bar a.cu会将-foo=bar传递给主机链接器。nvcc -forward-unknown-to-host-linker -foo bar a.cu将会对bar参数报错。nvcc -forward-unknown-to-host-linker -foo -bar a.cu会将-foo和-bar转发给主机链接器。
注意：在Windows系统上，还可参考选项-forward-slash-prefix-opts来处理以'/'开头的转发选项。
4.2.5.5. --forward-unknown-opts (-forward-unknown-opts)
意味着结合了选项 -forward-unknown-to-host-linker 和 -forward-unknown-to-host-compiler。
例如：
nvcc -forward-unknown-opts -foo=bar a.cu会将-foo=bar转发给主机链接器和编译器。nvcc -forward-unknown-opts -foo bar a.cu将会对bar参数报错。nvcc -forward-unknown-opts -foo -bar a.cu会将-foo和-bar转发给主机链接器和编译器。
注意：在Windows系统上，另请参阅选项-forward-slash-prefix-opts，用于转发以'/'开头的选项。
4.2.5.6. --forward-slash-prefix-opts (-forward-slash-prefix-opts)
如果指定此标志，并且启用了将未知选项转发给主机工具链（-forward-unknown-opts 或
-forward-unknown-to-host-linker 或 -forward-unknown-to-host-compiler），则以'/'开头的命令行参数将被转发给主机工具链。
例如：
nvcc -forward-slash-prefix-opts -forward-unknown-opts /T foo.cu会将标志/T转发给主机编译器和链接器。
当未指定此标志时，以'/'开头的命令行参数将被视为输入文件。
例如：
nvcc /T foo.cu会将'/T'视为输入文件，并使用Windows API函数GetFullPathName()来确定完整路径名。
注意：此标志仅在Windows系统上受支持。
4.2.5.7. --dont-use-profile (-noprof)
编译时不要使用nvcc.profile文件中的配置。
4.2.5.8. --threads number (-t)
指定用于并行执行编译步骤的最大线程数。
该选项可用于在针对多种架构进行编译时提升编译速度。编译器会创建number个线程来并行执行编译步骤。如果number为1，则忽略此选项。如果number为0，使用的线程数将等于机器上的CPU核心数。
4.2.5.9. --dryrun (-dryrun)
列出编译子命令但不执行它们。
4.2.5.10. --verbose (-v)
在执行编译子命令的同时列出它们。
4.2.5.11. --keep (-keep)
保留内部编译步骤中生成的所有中间文件。
4.2.5.12. --keep-dir directory (-keep-dir)
保留在此目录中生成的所有中间文件，这些文件是在内部编译步骤中产生的。
4.2.5.13. --save-temps (-save-temps)
该选项是--keep的别名。
4.2.5.14. --clean-targets (-clean)
删除所有非临时文件，这些文件是相同的nvcc命令在没有此选项时会生成的。
该选项会反转nvcc的行为。当指定此选项时，不会执行任何编译阶段。相反，nvcc原本会创建的所有非临时文件都将被删除。
4.2.5.15. --run-args arguments,... (-run-args)
当与--run一起使用时，为可执行文件指定命令行参数。
4.2.5.16. --use-local-env (-use-local-env)
使用此标志强制nvcc假设cl.exe的环境已设置完成，并跳过运行MSVC安装中用于设置cl.exe环境的批处理文件。这可以显著减少小型程序的整体编译时间。
4.2.5.17. --force-cl-env-setup (-force-cl-env-setup)
强制nvcc始终从MSVC安装中运行批处理文件，以设置cl.exe的环境（与旧版nvcc行为保持一致）。
如果未指定此标志，默认情况下，当满足以下条件时，nvcc将跳过批处理文件的执行：PATH中存在cl.exe、环境变量VSCMD_VER已设置，并且如果指定了-ccbin，则-ccbin指定的编译器需与PATH中的cl.exe匹配。跳过批处理文件执行可以显著减少小型程序的整体编译时间。
4.2.5.18. --input-drive-prefix prefix (-idp)
指定输入驱动前缀。
在Windows系统上，所有涉及文件名的命令行参数在传递给纯Windows可执行文件之前，都必须转换为Windows原生格式。此选项用于指定当前开发环境如何表示绝对路径。对于Cygwin构建环境，使用/cygwin/作为prefix；对于MinGW环境，则使用/作为prefix。
4.2.5.19. --dependency-drive-prefix prefix (-ddp)
指定依赖驱动前缀。
在Windows系统上，当生成依赖文件时（参见--generate-dependencies），所有文件名必须根据所使用的make实例进行适当转换。某些make实例在处理原生Windows格式的绝对路径中的冒号时会遇到问题，这取决于编译make实例的环境。对于Cygwin的make，使用/cygwin/作为prefix；对于MinGW，则使用/作为prefix。或者通过不指定任何内容，保留这些文件名的原生Windows格式。
4.2.5.20. --drive-prefix prefix (-dp)
指定驱动器前缀。
此选项将prefix同时指定为--input-drive-prefix和--dependency-drive-prefix。
4.2.5.21. --dependency-target-name target (-MT)
在生成依赖文件时指定生成规则的目标名称（参见--generate-dependencies）。
4.2.5.22. --no-align-double
指定在32位平台上不应将-malign-double作为编译器参数传递。
警告：这将导致ABI与CUDA内核ABI在某些64位类型上不兼容。
4.2.5.23. --no-device-link (-nodlink)
在链接对象文件时跳过设备链接步骤。
4.2.5.24. --allow-unsupported-compiler (-allow-unsupported-compiler)
禁用对受支持主机编译器版本的nvcc检查。
使用不受支持的主机编译器可能导致编译失败或运行时执行错误。风险自负。此选项在MacOS上无效。
4.2.6. CUDA编译的导向选项
4.2.6.1. --default-stream {legacy|null|per-thread} (-default-stream)
指定默认情况下，编译程序中的CUDA命令将发送到的流。
允许的值
legacy
CUDA传统流（每个上下文，隐式与其他流同步）
每线程
普通CUDA流（每个线程独立，不会与其他流隐式同步）
null
legacy的已弃用别名
默认
legacy 被用作默认流。
4.2.7. GPU代码生成的导向选项
4.2.7.1. --gpu-architecture (-arch)
指定必须为CUDA输入文件编译的NVIDIA虚拟GPU架构类别的名称。
除了下面简写形式的例外情况，使用此选项指定的架构必须是虚拟架构（例如compute_50）。通常，仅此选项不会触发为实际架构生成PTX的汇编（这是nvcc选项--gpu-code的作用，见下文）；相反，其目的是控制输入到PTX的预处理和编译。
为方便起见，在简单的nvcc编译场景中，支持以下简写形式。若未指定--gpu-code选项的值，则该选项默认采用--gpu-architecture的值。在此特殊情况下（作为上述描述的唯一例外），为--gpu-architecture指定的值可以是实际架构（如sm_50），此时nvcc会将指定的实际架构及其最接近的虚拟架构作为有效架构值。例如，nvcc --gpu-architecture=sm_50等价于nvcc --gpu-architecture=compute_50 --gpu-code=sm_50,compute_50。若指定了加速版实际GPU（如-arch=sm_90a），则加速与非加速虚拟代码都会被加入代码列表：--gpu-architecture=compute_90a --gpu-code=sm_90a,compute_90,compute_90a。
当指定-arch=native时，nvcc会检测系统中可见的GPU并为其生成代码，此选项不会生成PTX程序。如果系统中没有可见的受支持GPU，则会发出警告，并使用默认架构。
如果指定了-arch=all，nvcc会为所有支持的架构(sm_*)嵌入编译后的代码镜像，并为最高主虚拟架构生成PTX程序。对于-arch=all-major，nvcc会为所有支持的主版本(sm_*0)嵌入编译后的代码镜像，同时包含最早支持的版本，并为最高主虚拟架构添加PTX程序。
查看虚拟架构功能列表了解支持的虚拟架构列表，以及GPU功能列表了解支持的实际架构列表。
默认
sm_52 被用作默认值；首先生成针对 compute_52 的 PTX 代码，然后针对 sm_52 进行汇编和优化。
4.2.7.2. --gpu-code code,... (-code)
指定用于组装和优化PTX的NVIDIA GPU名称。
nvcc 会在生成的可执行文件中嵌入针对每个指定代码架构的编译代码映像，其中包含针对每个真实架构（如sm_50）的纯二进制加载映像，以及针对虚拟架构（如compute_50）的PTX代码。
在运行时，如果找不到适用于当前GPU的二进制加载映像，CUDA运行时系统会动态编译此类嵌入式PTX代码。
为选项--gpu-architecture和--gpu-code指定的架构可以是虚拟的也可以是实际的，但code架构必须与arch架构兼容。当使用--gpu-code选项时，--gpu-architecture选项的值必须是虚拟PTX架构。
例如，--gpu-architecture=compute_60与--gpu-code=sm_52不兼容，因为早期编译阶段会假定存在compute_60特性，而这些特性在sm_52上并不具备。
4.2.7.3. --generate-code specification (-gencode)
此选项提供了一种通用化的方式，用于替代--gpu-architecture=arch --gpu-code=code,...选项组合，以指定nvcc在代码生成方面的行为。
当使用之前的选项为相同的虚拟架构生成不同实际架构的PTX代码时，--generate-code选项允许为不同的虚拟架构生成多个PTX版本。实际上，--gpu-architecture=arch --gpu-code=code,...等同于--generate-code=arch=arch,code=code,...。
--generate-code 选项可以为不同的虚拟架构重复指定。
4.2.7.4. --relocatable-device-code {true|false} (-rdc)
启用或禁用可重定位设备代码的生成。
如果禁用，将生成可执行的设备代码。可重定位设备代码必须先链接才能执行。
允许的值
truefalse
默认
可重定位设备代码的生成已禁用。
4.2.7.5. --entries entry,... (-e)
指定必须为其生成代码的全局入口函数。
为所有入口函数生成PTX代码，但仅对选定的入口函数进行汇编。此选项的入口函数名称必须以修饰名形式指定。
默认
nvcc 为所有入口函数生成代码。
4.2.7.6. --maxrregcount amount (-maxrregcount)
指定GPU函数可以使用的最大寄存器数量。
在达到函数特定限制之前，较高的值通常会提升执行该函数的单个GPU线程性能。但由于线程寄存器是从GPU全局寄存器池中分配的，此选项值越高，也会降低最大线程块大小，从而减少线程并行量。因此，最佳的maxrregcount值是权衡折衷的结果。
如果值小于ABI所需的最小寄存器数量，编译器会将其提升至ABI的最低限制。
用户程序可能无法使用所有寄存器，因为部分寄存器被编译器保留。
默认
未设定最大值上限。
4.2.7.7. --use_fast_math (-use_fast_math)
利用快速数学库。
--use_fast_math 意味着 --ftz=true --prec-div=false --prec-sqrt=false --fmad=true。
4.2.7.8. --ftz {true|false} (-ftz)
控制单精度非规格化数的支持。
--ftz=true 将非规格化数值刷新为零，而 --ftz=false 保留非规格化数值。
--use_fast_math 意味着 --ftz=true。
允许的值
truefalse
默认
此选项设置为false且nvcc保留非规格化值。
4.2.7.9. --prec-div {true|false} (-prec-div)
该选项控制单精度浮点除法和倒数运算。
--prec-div=true 启用IEEE四舍五入模式，而 --prec-div=false 启用快速近似模式。
--use_fast_math 表示 --prec-div=false。
允许的值
truefalse
默认
此选项设置为true时，nvcc将启用IEEE四舍五入模式。
4.2.7.10. --prec-sqrt {true|false} (-prec-sqrt)
该选项控制单精度浮点平方根运算。
--prec-sqrt=true 启用IEEE四舍五入模式，而 --prec-sqrt=false 启用快速近似模式。
--use_fast_math 表示 --prec-sqrt=false。
允许的值
truefalse
默认
此选项设置为true时，nvcc将启用IEEE四舍五入模式。
4.2.7.11. --fmad {true|false} (-fmad)
此选项启用（禁用）将浮点乘法和加法/减法运算合并为浮点乘加运算（FMAD、FFMA或DFMA）。
--use_fast_math 表示启用 --fmad=true。
允许的值
truefalse
默认
此选项设置为true时，nvcc会启用将浮点乘法和加法/减法收缩为浮点乘加运算（FMAD、FFMA或DFMA）的功能。
4.2.7.12. --extra-device-vectorization (-extra-device-vectorization)
此选项启用更激进的设备代码向量化。
4.2.7.13. --compile-as-tools-patch (-astoolspatch)
为CUDA工具编译补丁代码。隐含–keep-device-functions选项。
只能与--ptx、--cubin或--fatbin一起使用。
不得与 -rdc=true 或 -ewp 同时使用。
某些PTX ISA功能在此编译模式下可能无法使用。
4.2.7.14. --keep-device-functions (-keep-device-functions)
在整个程序编译模式下，保留用户定义的外部链接__device__函数定义在生成的PTX中。
4.2.7.15. --jump-table-density percentage (-jtd)
指定switch语句中的分支密度百分比，并将其作为最小阈值，用于判断是否使用跳转表(brx.idx指令)来实现switch语句。
百分比范围从0到101（包含边界值）。
默认
此选项设置为101时，nvcc会禁用switch语句的跳转表生成。
4.2.8. 通用工具选项
4.2.8.1. --disable-warnings (-w)
禁止显示所有警告信息。
4.2.8.2. --source-in-ptx (-src-in-ptx)
在PTX中交错源数据。
仅可与--device-debug或--generate-line-info配合使用。
4.2.8.3. --restrict (-restrict)
断言所有内核指针参数均为restrict指针。
4.2.8.4. --Wno-deprecated-gpu-targets (-Wno-deprecated-gpu-targets)
抑制关于已弃用GPU目标架构的警告。
4.2.8.5. --Wno-deprecated-declarations (-Wno-deprecated-declarations)
抑制对已弃用实体使用的警告。
4.2.8.6. --Wreorder (-Wreorder)
当成员初始化器被重新排序时生成警告。
4.2.8.7. --Wdefault-stream-launch (-Wdefault-stream-launch)
当在<<<...>>>内核启动语法中未提供显式流参数时生成警告。
4.2.8.8. --Wmissing-launch-bounds (-Wmissing-launch-bounds)
当__global__函数没有显式的__launch_bounds__注解时生成警告。
4.2.8.9. --Wext-lambda-captures-this (-Wext-lambda-captures-this)
当扩展lambda隐式捕获this时生成警告。
4.2.8.10. --Werror kind,... (-Werror)
将指定类型的警告转换为错误。
以下是该选项接受的警告类型列表：
all-warnings
将所有警告视为错误。
cross-execution-space-call
对不支持的跨执行空间调用采取更严格的限制。编译器将对从
__host____device__调用__host__函数的情况生成错误而非警告。
reorder
当成员初始化器被重新排序时生成错误。
default-stream-launch
当在
<<<...>>>内核启动语法中未提供显式流参数时生成错误。
missing-launch-bounds
当
__global__函数没有显式的__launch_bounds__注解时生成警告。
ext-lambda-captures-this
当扩展lambda隐式捕获
this时生成错误。
deprecated-declarations
在使用已弃用的实体时生成错误。
4.2.8.11. --display-error-number (-err-no)
此选项显示由CUDA前端编译器（注意：非主机编译器）生成的任何消息的诊断编号。
4.2.8.12. --no-display-error-number (-no-err-no)
此选项禁用显示由CUDA前端编译器生成的任何消息的诊断编号（注意：不是主机编译器）。
4.2.8.13. --diag-error errNum,... (-diag-error)
针对CUDA前端编译器生成的指定诊断消息发出错误（注意：不影响主机编译器/预处理器生成的诊断）。
4.2.8.14. --diag-suppress errNum,... (-diag-suppress)
抑制由CUDA前端编译器生成的指定诊断消息（注意：不影响主机编译器/预处理器生成的诊断信息）。
4.2.8.15. --diag-warn errNum,... (-diag-warn)
针对CUDA前端编译器生成的指定诊断消息发出警告（注意：不影响主机编译器/预处理器生成的诊断信息）。
4.2.8.16. --resource-usage (-res-usage)
显示GPU代码的资源使用情况，例如寄存器和内存。
当设置--relocatable-device-code=true时，此选项隐含--nvlink-options=--verbose。否则，它隐含--ptxas-options=--verbose。
4.2.8.17. --device-stack-protector {true|false} (-device-stack-protector)
启用或禁用设备代码中堆栈保护机制的生成。
栈保护机制(Stack canaries)使得利用涉及栈局部变量的某些内存安全漏洞变得更加困难。编译器会使用启发式方法来评估每个函数中此类漏洞的风险。只有那些被判定为高风险的函数才会使用栈保护机制。
允许的值
truefalse
默认
设备代码中的堆栈保护机制生成已被禁用。
4.2.8.18. --help (-h)
打印此工具的帮助信息。
4.2.8.19. --version (-V)
打印此工具的版本信息。
4.2.8.20. --options-file file,... (-optf)
从指定文件包含命令行选项。
4.2.8.21. --time filename (-time)
生成一个逗号分隔值表格，记录每个编译阶段所花费的时间，并将其追加到作为选项参数给出的文件末尾。如果文件为空，表格的第一行将生成列标题。
如果文件名为-，计时数据将在标准输出中生成。
4.2.8.22. --qpp-config config (-qpp-config)
在使用q++主机编译器时指定配置（[[编译器/]版本,][目标]）。该参数将通过q++编译器的-V标志传递。
4.2.8.23. --list-gpu-code (-code-ls)
列出该工具支持的非加速GPU架构(sm_XX)并退出。
如果同时设置了–list-gpu-code和–list-gpu-arch，列表将使用与–generate-code值相同的格式显示。
4.2.8.24. --list-gpu-arch (-arch-ls)
列出该工具支持的非加速虚拟设备架构(compute_XX)并退出。
如果同时设置了–list-gpu-arch和–list-gpu-code，列表将使用与–generate-code值相同的格式显示。
4.2.9. 阶段选项
以下部分列出了一些对底层编译工具有用的选项。
4.2.9.1. Ptxas选项
下表列出了一些可通过nvcc选项-Xptxas指定的实用ptxas选项。
4.2.9.1.1. --allow-expensive-optimizations (-allow-expensive-optimizations)
启用（禁用）允许编译器使用最大可用资源（内存和编译时间）执行代价高昂的优化。
如果未指定，默认行为是为优化级别 >= O2 启用此功能。
4.2.9.1.2. --compile-only (-c)
生成可重定位对象。
4.2.9.1.3. --def-load-cache (-dlcm)
全局/通用加载的默认缓存修饰符。
4.2.9.1.4. --def-store-cache (-dscm)
全局/通用存储的默认缓存修饰符。
4.2.9.1.5. --device-debug (-g)
语义与 nvcc 选项 --device-debug 相同。
4.2.9.1.6. --disable-optimizer-constants (-disable-optimizer-consts)
禁用优化器常量库的使用。
4.2.9.1.7. --entry entry,... (-e)
语义与 nvcc 选项 --entries 相同。
4.2.9.1.8. --fmad (-fmad)
语义与 nvcc 选项 --fmad 相同。
4.2.9.1.9. --force-load-cache (-flcm)
在全局/通用加载上强制指定缓存修饰符。
4.2.9.1.10. --force-store-cache (-fscm)
在全局/通用存储上强制指定缓存修饰符。
4.2.9.1.11. --generate-line-info (-lineinfo)
语义与 nvcc 选项 --generate-line-info 相同。
4.2.9.1.12. --gpu-name gpuname (-arch)
指定NVIDIA GPU的名称以生成代码。
此选项也接受虚拟计算架构，在这种情况下会抑制代码生成。这可以仅用于解析。
Allowed values for this option: compute_50, compute_52, compute_53, compute_60,
compute_61, compute_62, compute_70, compute_72, compute_75, compute_80,
compute_86, compute_87, compute_89, compute_90, compute_90a, compute_100,
compute_100a, compute_101, compute_101a, compute_120, compute_120a, lto_50, lto_52, lto_53,
lto_60, lto_61, lto_62, lto_70, lto_72, lto_75, lto_80, lto_86,
lto_87, lto_89, lto_90, lto_90a, lto_100, lto_100a, lto_101, lto_101a,
sm_50, sm_52, sm_53, sm_60, sm_61, sm_62, sm_70, sm_72, sm_75,
sm_80, sm_86, sm_87, sm_89, sm_90, sm_90a, sm_100, sm_100a, sm_101,
sm_101a, sm_120, sm_120a
默认值: sm_52.
4.2.9.1.13. --help (-h)
语义与 nvcc 选项 --help 相同。
4.2.9.1.14. --machine (-m)
语义与 nvcc 选项 --machine 相同。
4.2.9.1.15. --maxrregcount amount (-maxrregcount)
语义与 nvcc 选项 --maxrregcount 相同。
4.2.9.1.16. --opt-level N (-O)
指定优化级别。
默认值: 3.
4.2.9.1.17. --options-file file,... (-optf)
语义与 nvcc 选项 --options-file 相同。
4.2.9.1.18. --position-independent-code (-pic)
生成位置无关代码。
默认值：
针对整个程序编译：true
否则: false
4.2.9.1.19. --preserve-relocs (-preserve-relocs)
此选项将使ptxas为变量生成可重定位引用，并在链接的可执行文件中保留为其生成的重定位信息。
4.2.9.1.20. --sp-bounds-check (-sp-bounds-check)
生成栈指针边界检查代码序列。
当指定--device-debug或--opt-level=0时，此选项会自动启用。
4.2.9.1.21. --suppress-async-bulk-multicast-advisory-warning (-suppress-async-bulk-multicast-advisory-warning)
在使用cp.async.bulk{.tensor}指令配合sm_90时，抑制关于.multicast::cluster修饰符使用的警告。
4.2.9.1.22. --verbose (-v)
启用详细模式，打印代码生成统计信息。
4.2.9.1.23. --version (-V)
语义与 nvcc 选项 --version 相同。
4.2.9.1.24. --warning-as-error (-Werror)
将所有警告视为错误。
4.2.9.1.25. --warn-on-double-precision-use (-warn-double-usage)
警告：如果在指令中使用了双精度浮点数。
4.2.9.1.26. --warn-on-local-memory-usage (-warn-lmem-usage)
如果使用了本地内存则发出警告。
4.2.9.1.27. --warn-on-spills (-warn-spills)
如果寄存器溢出到本地内存则发出警告。
4.2.9.1.28. --compile-as-tools-patch (-astoolspatch)
为CUDA工具编译补丁代码。
不得与 -Xptxas -c 或 -ewp 同时使用。
某些PTX ISA功能在此编译模式下可能无法使用。
4.2.9.1.29. --maxntid (-maxntid)
指定一个线程块可以拥有的最大线程数。
如果与-maxrregcount选项同时使用，此选项将被忽略。对于已指定.maxntid指令的入口函数，此选项同样会被忽略。
4.2.9.1.30. --minnctapersm (-minnctapersm)
指定要映射到SM的最小CTA数量。
如果与-maxrregcount选项同时使用，此选项将被忽略。对于指定了.minnctapersm指令的入口函数，此选项同样会被忽略。
4.2.9.1.31. --override-directive-values (-override-directive-values)
通过对应的选项值覆盖PTX指令值。
此选项仅对-minnctapersm、-maxntid和-maxrregcount选项有效。
4.2.9.1.32. --make-errors-visible-at-exit (-make-errors-visible-at-exit)
在退出点生成必要的指令，使内存故障和错误在退出时可见。
4.2.9.1.33. --oFast-compile (-Ofc)
指定级别以优先考虑设备代码的编译速度。
默认值: 0.
4.2.9.1.34. --device-stack-protector (-device-stack-protector)
启用或禁用设备代码中堆栈保护机制的生成。
栈保护机制(Stack canaries)使得利用涉及栈局部变量的某些内存安全漏洞变得更加困难。编译器会使用启发式方法来评估每个函数中此类漏洞的风险。只有那些被判定为高风险的函数才会使用栈保护机制。
允许的值
truefalse
默认
设备代码中的堆栈保护机制生成已被禁用。
4.2.9.1.35. --g-tensor-memory-access-check (-g-tmem-access-check)
为tcgen05操作启用张量内存访问检查。
4.2.9.1.36. --split-compile (-split-compile)
指定运行编译器优化时使用的最大并发线程数。
如果指定的值为1，该选项将被忽略。
如果指定的值为0，则线程数将等于底层机器的CPU数量。
4.2.9.2. NVLINK 选项
以下是一些有用的nvlink选项列表，可以通过nvcc的--nvlink-options选项指定。
4.2.9.2.1. --disable-warnings (-w)
禁止显示所有警告信息。
4.2.9.2.2. --preserve-relocs (-preserve-relocs)
在链接的可执行文件中保留已解析的重定位信息。
4.2.9.2.3. --verbose (-v)
启用详细模式，打印代码生成统计信息。
4.2.9.2.4. --warning-as-error (-Werror)
将所有警告视为错误。
4.2.9.2.5. --suppress-arch-warning (-suppress-arch-warning)
抑制当对象不包含目标架构代码时通常会打印的警告。
4.2.9.2.6. --suppress-stack-size-warning (-suppress-stack-size-warning)
抑制当无法确定堆栈大小时通常会打印的警告信息。
4.2.9.2.7. --dump-callgraph (-dump-callgraph)
转储有关调用图和寄存器使用情况的信息。
4.2.9.2.8. --dump-callgraph-no-demangle (-dump-callgraph-no-demangle)
转储调用图信息而不进行名称还原。
4.2.9.2.9. --Xptxas (-Xptxas)
Ptxas选项（仅在LTO模式下使用）。
4.2.9.2.10. --cpu-arch (-cpu-arch)
指定CPU目标架构的名称。
4.2.9.2.11. --extra-warnings (-extrawarn)
针对可能出现的问题发出额外警告。
4.2.9.2.12. --gen-host-linker-script (-ghls)
指定要生成的主机链接器脚本类型。
4.2.9.2.13. --ignore-host-info (-ignore-host-info)
忽略有关主机引用的信息，因此不要移除可能被主机引用的设备代码。
4.2.9.2.14. --keep-system-libraries (-keep-system-libraries)
不要优化掉系统库（例如cudadevrt）代码。
4.2.9.2.15. --kernels-used (-kernels-used)
指定所使用的内核。可以是内核名称的一部分，因此名称中包含该字符串的任何内核都会被匹配。如果使用此选项，则其他所有内核将被视为死代码并被移除。
4.2.9.2.16. --options-file (-optf)
从指定文件中包含命令行选项。
4.2.9.2.17. --report-arch (-report-arch)
在错误消息中报告SM目标架构。
4.2.9.2.18. --suppress-debug-info (-suppress-debug-info)
不在输出中保留调试符号。如果未与–debug选项一起使用，则忽略此选项。
4.2.9.2.19. --variables-used (-variables used)
指定所使用的变量。可以是变量名的一部分，因此名称中包含该字符串的任何变量都会被匹配。如果使用此选项，则其他任何变量都将被视为死代码，除非设备代码中有其他访问，否则可能会被移除。
4.2.9.2.20. --device-stack-protector {true|false} (-device-stack-protector)
启用或禁用设备代码中堆栈保护机制的生成（仅在使用LTO时有效）。
栈保护机制(Stack canaries)使得利用涉及栈局部变量的某些内存安全漏洞变得更加困难。编译器会使用启发式方法来评估每个函数中此类漏洞的风险。只有那些被判定为高风险的函数才会使用栈保护机制。
允许的值
truefalse
默认
设备代码中的堆栈保护机制生成已被禁用。
4.3. NVCC环境变量
NVCC_PREPEND_FLAGS 和 NVCC_APPEND_FLAGS：
可以通过设置以下环境变量来增强nvcc命令行标志的功能：
NVCC_PREPEND_FLAGS
在正常nvcc命令行之前需要注入的标志。
NVCC_APPEND_FLAGS
在正常nvcc命令行后需要注入的标志。
例如，设置后：
export NVCC_PREPEND_FLAGS='-G -keep -arch=sm_60'
export NVCC_APPEND_FLAGS='-DNAME=" foo "'
以下调用方式：
nvcc foo.cu -o foo
等同于：
nvcc -G -keep -arch=sm_60 foo.cu -o foo -DNAME=" foo "
这些环境变量可以用于全局注入nvcc标志，而无需修改构建脚本。
来自NVCC_PREPEND_FLAGS或NVCC_APPEND_FLAGS的附加标志将在详细日志(--verbose)中列出。
NVCC_CCBIN:
可以使用环境变量NVCC_CCBIN来设置默认的主机编译器。例如，设置后：
export NVCC_CCBIN='gcc'
nvcc 会在未设置 --compiler-bindir 时选择 gcc 作为主机编译器。
NVCC_CCBIN 可用于全局控制默认主机编译器。如果同时设置了 NVCC_CCBIN 和 --compiler-bindir，nvcc 将优先选择由 --compiler-bindir 指定的主机编译器。例如：
export NVCC_CCBIN='gcc'
nvcc foo.cu -ccbin='clang' -o foo
在这种情况下，nvcc会选择clang作为主机编译器。
5. GPU编译
本章描述了由nvcc与CUDA驱动程序共同维护的GPU编译模型。内容包含若干技术章节，并在最后提供了具体示例。
5.1. GPU 世代
为了支持架构演进，NVIDIA GPU按不同代次发布。新一代产品会在功能和/或芯片架构上带来重大改进，而同代GPU型号之间则存在细微的配置差异，这些差异会适度影响功能、性能或两者兼具。
GPU应用程序的二进制兼容性在不同代际之间无法保证。例如，为Fermi GPU编译的CUDA应用程序很可能无法在Kepler GPU上运行（反之亦然）。这是因为每一代的指令集和指令编码与其他代际不同。
在同一代GPU中，在特定条件下可以保证二进制兼容性，因为它们共享基本指令集。这种情况适用于两个GPU版本不存在功能差异（例如一个版本是另一个版本的缩减版），或者一个版本的功能完全包含于另一个版本中。后者的一个例子是基础Maxwell版本sm_50，其功能是所有其他Maxwell版本的子集：任何为sm_50编译的代码都可以在所有其他Maxwell GPU上运行。
5.2. GPU功能列表
下表列出了当前GPU架构的名称，并标注了它们提供的功能特性。还存在其他差异，如寄存器数量和处理器集群规模等，这些仅影响执行性能。
在CUDA命名方案中，GPU被命名为sm_xy，其中x表示GPU代际编号，y表示该代际中的版本号。此外，为了便于比较GPU功能，CUDA尝试选择这样的GPU命名方式：如果x1y1 <= x2y2，那么sm_x1y1的所有非ISA相关功能都包含在sm_x2y2中。由此可以得出sm_50确实是Maxwell架构的基础型号，这也解释了为什么表格中更高版本的条目总是对较低版本的功能扩展。此外，如果我们从指令编码的角度抽象来看，这意味着sm_50的功能将持续包含在所有后续GPU代际中。正如我们接下来将看到的，这一特性将成为nvcc实现应用程序兼容性支持的基础。
|
Maxwell支持 |
|
Pascal架构支持 |
|
Volta架构支持 |
|
图灵支持 |
|
支持NVIDIA Ampere GPU架构 |
|
Ada支持 |
|
Hopper支持 |
|
Blackwell 支持 |
5.3. 应用程序兼容性
二进制代码在CPU代际间的兼容性，加上公开的指令集架构，是确保分布式应用在现网环境中能够持续运行于新一代主流CPU的常规机制。
对于GPU来说情况有所不同，因为NVIDIA无法在不牺牲常规GPU改进机会的情况下保证二进制兼容性。相反，正如图形编程领域已经形成的惯例，nvcc采用两阶段编译模型来确保应用程序与未来GPU世代的兼容性。
5.4. 虚拟架构
GPU编译通过一种称为PTX的中间表示进行，可以将其视为虚拟GPU架构的汇编语言。与实际图形处理器不同，这种虚拟GPU完全由其向应用程序提供的能力或功能集定义。具体而言，虚拟GPU架构提供（很大程度上）通用的指令集，且二进制指令编码不是问题，因为PTX程序始终以文本格式表示。
因此，nvcc编译命令总是使用两种架构：一个虚拟中间架构，加上一个真实GPU架构来指定目标执行处理器。要使这样的nvcc命令有效，真实架构必须是虚拟架构的实现。下文将对此进行进一步说明。
所选的虚拟架构更多是声明应用程序所需的GPU能力：使用较小的虚拟架构仍允许在第二个nvcc阶段选择更广泛的实际架构范围。相反，如果指定了应用程序未使用的虚拟架构功能，则会不必要地限制在第二个nvcc阶段可以指定的可能GPU集合。
由此可见，虚拟架构应尽可能选择较低的版本，从而最大化实际可运行的GPU数量。而真实架构则应尽可能选择较高的版本（假设更高版本总能生成更优代码），但这需要预先了解应用程序实际将运行在哪些GPU上。正如我们稍后在即时编译场景中将看到的——当驱动程序确切知晓运行时GPU信息时，程序即将启动/执行的GPU就是目标运行设备。
采用虚拟与真实架构的两阶段编译
5.5. 虚拟架构功能列表
|
Maxwell支持 |
|
Pascal架构支持 |
|
Volta架构支持 |
|
图灵支持 |
|
支持NVIDIA Ampere GPU架构 |
|
Ada支持 |
|
Hopper支持 |
|
Blackwell 支持 |
上表列出了当前定义的虚拟架构。虚拟架构的命名规则与GPU功能列表章节中展示的真实架构命名方案相同。
5.6. 更多机制
显然，编译阶段本身并不能直接实现应用程序与未来GPU的兼容性目标。为此，我们还需要另外两种机制：即时编译（JIT）和胖二进制文件。
5.6.1. 即时编译
将代码编译到实际GPU的过程会将其绑定到某一代GPU硬件。在该代GPU范围内，需要在覆盖范围和潜在性能之间做出权衡。例如，编译为sm_50可使代码在所有Maxwell架构GPU上运行，但如果目标仅限于Maxwell GM206及后续型号，编译为sm_53可能会生成更优化的代码。
设备代码的即时编译
通过指定一个虚拟代码架构而非真实GPU，nvcc将PTX代码的汇编推迟到应用程序运行时，此时目标GPU已完全确定。例如，当应用程序在sm_50或更高架构上启动时，以下命令允许生成完全匹配的GPU二进制代码。
nvcc x.cu --gpu-architecture=compute_50 --gpu-code=compute_50
即时编译的缺点是会增加应用程序启动延迟，但可以通过让CUDA驱动程序使用编译缓存来缓解这一问题（参见CUDA C++编程指南的"3.1.1.2. 即时编译"章节），该缓存在应用程序多次运行期间是持久化的。
5.6.2. Fatbinaries
另一种解决方案是通过指定多个代码实例来克服JIT启动延迟，同时仍允许在新款GPU上执行，如下所示
nvcc x.cu --gpu-architecture=compute_50 --gpu-code=compute_50,sm_50,sm_52
该命令为两种Maxwell架构变体生成精确代码，同时提供PTX代码以便在遇到下一代GPU时供JIT使用。nvcc将其设备代码组织为fatbinaries（胖二进制文件），这种格式能够保存同一GPU源代码的多种转换版本。在运行时，当启动设备函数时，CUDA驱动程序将选择最合适的转换版本。
5.7. NVCC 示例
5.7.1. 基础表示法
nvcc provides the options --gpu-architecture and --gpu-code for specifying the target architectures for both translation stages. Except for allowed short hands described below, the --gpu-architecture option takes a single value, which must be the name of a virtual compute architecture, while option --gpu-code takes a list of values which must all be the names of actual GPUs. nvcc performs a stage 2 translation for each of these GPUs, and will embed the result in the result of compilation (which usually is a host object file or executable).
示例
nvcc x.cu --gpu-architecture=compute_50 --gpu-code=sm_50,sm_52
5.7.2. 简写形式
nvcc 为简单场景提供了一些简写方式。
5.7.2.1. 简写形式1
--gpu-code arguments can be virtual architectures. In this case the stage 2 translation will be omitted for such virtual architecture, and the stage 1 PTX result will be embedded instead. At application launch, and in case the driver does not find a better alternative, the stage 2 compilation will be invoked by the driver with the PTX as input.
示例
nvcc x.cu --gpu-architecture=compute_50 --gpu-code=compute_50,sm_50,sm_52
5.7.2.2. 简写形式2
--gpu-code选项可以省略。仅在这种情况下，--gpu-architecture的值可以是非虚拟架构。--gpu-code值默认采用由--gpu-architecture指定的GPU所实现的最接近虚拟架构，再加上--gpu-architecture值本身。这个最接近的虚拟架构将作为有效的--gpu-architecture值使用。如果--gpu-architecture值是虚拟架构，它也会被用作有效的--gpu-code值。
示例
nvcc x.cu --gpu-architecture=sm_52
nvcc x.cu --gpu-architecture=compute_50
等同于
nvcc x.cu --gpu-architecture=compute_52 --gpu-code=sm_52,compute_52
nvcc x.cu --gpu-architecture=compute_50 --gpu-code=compute_50
5.7.2.3. 简写3
--gpu-architecture 和 --gpu-code 选项都可以省略。
示例
nvcc x.cu
等同于
nvcc x.cu --gpu-architecture=compute_52 --gpu-code=sm_52,compute_52
5.7.3. 扩展表示法
选项 --gpu-architecture 和 --gpu-code 可用于所有需要为一个或多个GPU生成代码的情况，使用共同的虚拟架构。这将导致单次调用 nvcc 阶段1（即预处理和生成虚拟PTX汇编代码），然后为每个指定的GPU重复阶段2（二进制代码生成）的编译过程。
使用统一的虚拟架构意味着在整个nvcc编译过程中，所有预设的GPU特性都是固定的。例如，以下nvcc命令同时为sm_50代码和sm_53代码假设不支持半精度浮点运算：
nvcc x.cu --gpu-architecture=compute_50 --gpu-code=compute_50,sm_50,sm_53
有时需要针对不同架构执行不同的GPU代码生成步骤。这可以通过使用nvcc选项--generate-code来实现，此时必须用它来替代--gpu-architecture和--gpu-code的组合。
与选项--gpu-architecture不同，选项--generate-code可以在nvcc命令行中重复使用。它接受子选项arch和code，这些子选项不能与它们对应的主选项混淆，但行为类似。如果使用重复架构编译，则设备代码必须基于架构标识宏__CUDA_ARCH__的值进行条件编译，该宏将在下一节中描述。
例如，以下假设sm_50和sm_52架构不支持半精度浮点运算，但sm_53架构完全支持：
nvcc x.cu \
--generate-code arch=compute_50,code=sm_50 \
--generate-code arch=compute_50,code=sm_52 \
--generate-code arch=compute_53,code=sm_53
或者，将实际的GPU代码生成工作交给CUDA驱动中的JIT编译器处理：
nvcc x.cu \
--generate-code arch=compute_50,code=compute_50 \
--generate-code arch=compute_53,code=compute_53
代码子选项可以通过稍微复杂一些的语法组合使用：
nvcc x.cu \
--generate-code arch=compute_50,code=[sm_50,sm_52] \
--generate-code arch=compute_53,code=sm_53
5.7.4. 虚拟架构宏
架构标识宏__CUDA_ARCH__会被赋予一个三位数字符串值xy0(以字面量0结尾)，用于每个阶段1的nvcc编译，该编译针对compute_xy进行。
该宏可用于GPU函数的实现中，用于确定当前正在编译的虚拟架构。主机代码（非GPU代码）不得依赖于此宏。
架构列表宏__CUDA_ARCH_LIST__是一个逗号分隔的__CUDA_ARCH__值列表，对应编译器调用中指定的每个虚拟架构。该列表按数值升序排列。
宏定义 __CUDA_ARCH_LIST__ 在编译C、C++和CUDA源文件时会被定义。
例如，以下nvcc编译命令行会将__CUDA_ARCH_LIST__定义为500,530,800：
nvcc x.cu \
--generate-code arch=compute_80,code=sm_80 \
--generate-code arch=compute_50,code=sm_52 \
--generate-code arch=compute_50,code=sm_50 \
--generate-code arch=compute_53,code=sm_53
6. 在CUDA中使用独立编译
在5.0版本发布之前，CUDA不支持单独编译，因此CUDA代码无法跨文件调用设备函数或访问变量。这种编译方式被称为整体程序编译。我们一直支持主机代码的单独编译，只是CUDA设备代码需要全部集中在一个文件中。从CUDA 5.0开始，支持设备代码的单独编译，但旧的整程序模式仍然是默认设置，因此新增了调用单独编译的选项。
6.1. 独立编译的代码修改
为实现设备代码的单独编译所需的代码修改与您已为主机代码所做的相同，即使用extern和static来控制符号的可见性。请注意，以前在CUDA代码中extern是被忽略的；现在它将被遵循。通过使用static，可以在不同文件中拥有多个同名的设备符号。因此，通过字符串名称引用符号的CUDA API调用已被弃用；应改为通过符号地址来引用。
6.2. NVCC独立编译选项
CUDA的工作原理是将设备代码嵌入到主机对象中。在全程序编译模式下，它将可执行的设备代码嵌入到主机对象内。而在分离编译模式下，我们会将可重定位的设备代码嵌入主机对象，并运行设备链接器nvlink将所有设备代码链接在一起。随后，nvlink的输出会与所有主机对象通过主机链接器进行链接，最终生成可执行文件。
可重定位设备代码与可执行设备代码的生成由--relocatable-device-code选项控制。
--compile选项已用于控制在主机对象处停止编译，因此新增了--device-c选项，该选项仅执行--relocatable-device-code=true --compile操作。
要仅调用设备链接器，可以使用--device-link选项，该选项会生成一个包含嵌入式可执行设备代码的主机对象文件。然后必须将该输出传递给主机链接器。或者：
nvcc <objects>
可用于隐式调用设备和主机链接器。这是因为如果设备链接器没有看到任何可重定位代码，它就不会执行任何操作。
下图展示了流程。
CUDA 独立编译轨迹
6.3. 库
设备链接器能够读取静态主机库格式（Linux和Mac OS X上的.a，Windows上的.lib）。它会忽略任何动态库（.so或.dll）。--library和--library-path选项可用于向设备和主机链接器传递库文件。当使用--library选项时，库名称需不带库文件扩展名。
nvcc --gpu-architecture=sm_50 a.o b.o --library-path=<path> --library=foo
或者，在Windows系统上，可以不使用--library选项，直接指定包含库文件扩展名的库名称。
nvcc --gpu-architecture=sm_50 a.obj b.obj foo.lib --library-path=<path>
请注意，设备链接器会忽略任何不包含可重定位设备代码的对象。
6.4. 示例
假设我们有以下文件：
//---------- b.h ----------
#define N 8
extern __device__ int g[N];
extern __device__ void bar(void);
//---------- b.cu ----------
#include "b.h"
__device__ int g[N];
__device__ void bar (void)
{
g[threadIdx.x]++;
}
//---------- a.cu ----------
#include <stdio.h>
#include "b.h"
__global__ void foo (void) {
__shared__ int a[N];
a[threadIdx.x] = threadIdx.x;
__syncthreads();
g[threadIdx.x] = a[blockDim.x - threadIdx.x - 1];
bar();
}
int main (void) {
unsigned int i;
int *dg, hg[N];
int sum = 0;
foo<<<1, N>>>();
if(cudaGetSymbolAddress((void**)&dg, g)){
printf("couldn't get the symbol addr\n");
return 1;
}
if(cudaMemcpy(hg, dg, N * sizeof(int), cudaMemcpyDeviceToHost)){
printf("couldn't memcpy\n");
return 1;
}
for (i = 0; i < N; i++) {
sum += hg[i];
}
if (sum == 36) {
printf("PASSED\n");
} else {
printf("FAILED (%d)\n", sum);
}
return 0;
}
可以使用以下命令进行编译（这些示例适用于Linux系统）：
nvcc --gpu-architecture=sm_50 --device-c a.cu b.cu
nvcc --gpu-architecture=sm_50 a.o b.o
如果想分别调用设备和主机链接器，可以执行以下操作：
nvcc --gpu-architecture=sm_50 --device-c a.cu b.cu
nvcc --gpu-architecture=sm_50 --device-link a.o b.o --output-file link.o
g++ a.o b.o link.o --library-path=<path> --library=cudart
请注意，所有目标架构都必须传递给设备链接器，因为这决定了最终可执行文件中的内容（某些对象或库可能包含针对多个架构的设备代码，链接步骤随后可以选择将哪些内容放入最终可执行文件中）。
如果你想使用驱动API加载一个链接的cubin文件，可以仅请求cubin部分：
nvcc --gpu-architecture=sm_50 --device-link a.o b.o \
--cubin --output-file link.cubin
这些对象可以放入库中并与以下方式一起使用：
nvcc --gpu-architecture=sm_50 --device-c a.cu b.cu
nvcc --lib a.o b.o --output-file test.a
nvcc --gpu-architecture=sm_50 test.a
请注意，设备链接器仅支持静态库。
PTX文件可以编译为主机目标文件，然后通过以下方式链接：
nvcc --gpu-architecture=sm_50 --device-c a.ptx
一个使用库、主机链接器和动态并行功能的示例如下：
nvcc --gpu-architecture=sm_50 --device-c a.cu b.cu
nvcc --gpu-architecture=sm_50 --device-link a.o b.o --output-file link.o
nvcc --lib --output-file libgpu.a a.o b.o link.o
g++ host.o --library=gpu --library-path=<path> \
--library=cudadevrt --library=cudart
可以在单个主机可执行文件中进行多个设备链接，只要每个设备链接彼此独立。这种独立性要求意味着它们不能在设备可执行文件之间共享代码，也不能共享地址（例如，只有当设备链接同时看到调用方和潜在回调被调用方时，才能将设备函数地址从主机传递到设备进行回调；不能将地址从一个设备可执行文件传递到另一个设备可执行文件，因为它们是独立的地址空间）。
6.5. 独立编译优化
单独编译的代码可能无法达到整个程序代码的高性能，因为无法跨文件内联代码。一种仍能获得最佳性能的方法是使用链接时优化，该技术存储中间代码，然后链接在一起执行高级优化。这可以通过--dlink-time-opt或-dlto选项实现。该选项必须在编译和链接时同时指定。如果只有部分文件使用-dlto编译，那么这些文件将被链接并一起优化，而其余文件则使用正常的单独编译。副作用是将部分编译时间转移到链接阶段，对于非常大的代码可能存在一些可扩展性问题。如果要使用-gencode为多种架构编译，请使用-dc -gencode arch=compute_NN,code=lto_NN指定要存储的中间IR（其中NN是SM架构版本）。然后使用-dlto选项为特定架构进行链接。
从CUDA 12.0开始，通过nvJitLink库支持运行时LTO。
6.6. 潜在的独立编译问题
6.6.1. 对象兼容性
只有具有相同ABI版本、链接兼容的SM目标架构以及相同指针大小（32位或64位）的可重定位设备代码才能链接在一起。链接器的工具包版本必须≥对象的工具包版本。不兼容的对象将产生链接错误。链接兼容的SM架构是指具有兼容SASS二进制文件、无需转换即可组合的架构，例如sm_52和sm_50。对象可能已针对不同架构编译，但也可能提供PTX，在这种情况下，设备链接器会将PTX即时编译为所需架构的cubin，然后进行链接。可重定位设备代码需要CUDA 5.0或更高版本的工具包。
如果使用链接时间优化（Link Time Optimization）并配合-dlto参数，中间生成的LTOIR文件仅保证在主版本号内兼容（例如12.0和12.1版本的LTO中间文件可以链接，但12.1和11.6版本则不行）。
如果内核通过launch_bounds属性或--maxrregcount选项限制了寄存器数量，那么该内核调用的所有函数使用的寄存器数量不得超过该限制；如果超出限制，将会产生链接错误。
6.6.2. JIT链接支持
JIT链接指的是在加载时对代码进行隐式重新链接。如果在加载时cubin与目标架构不匹配，驱动程序会重新调用设备链接器为目标架构生成cubin，具体方式是先将每个对象的PTX通过JIT编译为对应的cubin，然后将新的cubin链接在一起。如果某个对象找不到针对目标架构的PTX或cubin，则链接将失败。目前不支持对LTO中间文件进行隐式JIT链接，但可以使用nvJitLink库显式链接它们。
6.6.3. 隐式CUDA主机代码
像上面b.cu这样的文件仅包含CUDA设备代码，因此有人可能认为不需要将b.o对象文件传递给主机链接器。但实际上，每当可以通过主机端访问设备符号时（无论是通过启动还是像cudaGetSymbolAddress()这样的API调用），都会生成隐式主机代码。这些隐式主机代码被放入b.o中，并且需要传递给主机链接器。此外，为了使JIT链接正常工作，所有设备代码都必须传递给主机链接器，否则主机可执行文件将不包含JIT链接所需的设备代码。因此，一个通用规则是：设备链接器和主机链接器必须看到相同的主机对象文件（如果这些对象文件中包含任何设备引用——如果文件是纯主机的，则设备链接器不需要看到它）。如果未将包含设备代码的对象文件传递给主机链接器，则会看到关于函数__cudaRegisterLinkedBinary_name调用未定义或未解析符号__fatbinwrap_name的错误消息。
6.6.4. 使用__CUDA_ARCH__
在单独编译时，__CUDA_ARCH__不能在头文件中使用，以避免不同对象可能包含不同行为。或者，必须确保所有对象都将为相同的compute_arch编译。如果在头文件中定义了一个弱函数或模板函数，并且其行为依赖于__CUDA_ARCH__，那么当对象为不同的计算架构编译时，这些对象中的函数实例可能会发生冲突。例如，如果a.h包含：
template<typename T>
__device__ T* getptr(void)
{
#if __CUDA_ARCH__ == 500
return NULL; /* no address */
#else
__shared__ T arr[256];
return arr;
#endif
}
如果a.cu和b.cu都包含了a.h并为相同类型实例化了getptr，且b.cu期望获取非NULL地址，则使用以下命令编译：
nvcc --gpu-architecture=compute_50 --device-c a.cu
nvcc --gpu-architecture=compute_52 --device-c b.cu
nvcc --gpu-architecture=sm_52 a.o b.o
在链接时只会使用一个版本的getptr，因此行为将取决于选择了哪个版本。为避免这种情况，必须为相同的计算架构编译a.cu和b.cu，或者在共享头文件函数中不应使用__CUDA_ARCH__。
6.6.5. 库中的设备代码
如果一个具有非弱外部链接的设备函数同时在库和非库对象（或另一个库）中定义，设备链接器会报多重定义错误（这与传统主机链接器不同，后者如果在较早的对象中已找到该函数定义，可能会忽略库对象中的函数定义）。
7. NVCC 其他用法
7.1. 跨平台编译
交叉编译通过以下nvcc命令行选项控制：
--compiler-bindir用于交叉编译场景，当底层主机编译器能够为目标平台生成对象时使用。
在x86系统上，如果CUDA工具包安装已配置为同时支持Tegra和非Tegra ARM目标的交叉编译，那么当指定了ARM主机交叉编译器时，nvcc将默认使用非Tegra配置。若要改用Tegra配置，请向nvcc传递"-target-dir aarch64-linux"参数。
7.2. 保留中间阶段文件
nvcc 默认将中间结果存储在临时文件中，这些文件会在程序完成前立即删除。根据当前操作系统的不同，所使用的临时文件目录位置如下：
- Windows
将使用环境变量
TEMP的值。如果未设置该变量，则改用C:\Windows\temp。- Other Platforms
将使用环境变量
TMPDIR的值。如果未设置该变量，则改用/tmp。
选项 --keep 会让 nvcc 将这些中间文件保存在当前目录或由 --keep-dir 指定的目录中，文件名规则如 Supported Phases 所述。
7.3. 清理生成的文件
由特定nvcc命令生成的所有文件可以通过重复该命令来清理，但需要添加--clean-targets选项。该选项在使用--keep后特别有用，因为--keep选项通常会留下大量中间文件。
由于使用--clean-targets会精确删除原始nvcc命令生成的内容，因此必须完全重复原始命令中的所有选项。例如，在以下示例中，省略--keep或添加--compile会产生不同的清理效果。
nvcc acos.cu --keep
nvcc acos.cu --keep --clean-targets
7.4. 打印代码生成统计信息
通过向nvcc传递--resource-usage选项，可以打印出每个已编译设备函数使用的寄存器数量和所需内存量的摘要信息。
$ nvcc --resource-usage acos.cu -arch sm_80
ptxas info : 1536 bytes gmem
ptxas info : Compiling entry function 'acos_main' for 'sm_80'
ptxas info : Function properties for acos_main
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 6 registers, 1536 bytes smem, 32 bytes cmem[0]
如上例所示，列出了静态分配的全局内存(gmem)的大小。
全局内存和部分常量存储区是模块作用域的资源，而非每个内核的资源。将常量变量分配到常量存储区是特定于配置文件的。
随后，会打印每个内核的资源信息。
堆栈帧是该函数使用的每线程堆栈占用情况。溢出存储和加载表示在堆栈内存上进行的存储和加载操作，这些操作用于存储无法分配到物理寄存器的变量。
同样会显示寄存器数量、共享内存大小以及常量存储区分配的总空间。
8. 通知
8.1. 通知
本文档仅供信息参考之用，不应视为对产品功能、状态或质量的保证。NVIDIA公司（“NVIDIA”）对本文件所含信息的准确性或完整性不作任何明示或暗示的陈述或保证，并对其中可能存在的错误不承担任何责任。NVIDIA对于因使用此类信息而产生的后果、或因使用该信息导致的第三方专利或其他权利侵权概不负责。本文件不构成对开发、发布或交付任何材料（定义见下文）、代码或功能的承诺。
NVIDIA保留随时对本文件进行更正、修改、增强、改进以及任何其他变更的权利，恕不另行通知。
客户在下单前应获取最新的相关信息，并确认这些信息是最新且完整的。
除非NVIDIA与客户授权代表签署的单独销售协议中另有约定，否则NVIDIA产品的销售均以订单确认时提供的NVIDIA标准销售条款和条件为准（以下简称"销售条款"）。NVIDIA特此明确反对将任何客户通用条款适用于本文件所述NVIDIA产品的采购。本文件不直接或间接构成任何合同义务。
NVIDIA产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命支持设备，也不适用于那些可以合理预期NVIDIA产品故障或失灵会导致人身伤害、死亡、财产或环境损害的应用场景。NVIDIA对于在此类设备或应用中使用和/或包含NVIDIA产品不承担任何责任，因此客户需自行承担相关风险。
NVIDIA不声明或保证基于本文档的产品适用于任何特定用途。NVIDIA未必会对每个产品的所有参数进行测试。客户应全权负责评估和确定本文档所含信息的适用性，确保产品适合并满足客户计划的应用需求，并执行必要的应用测试以避免应用或产品出现故障。客户产品设计中的缺陷可能会影响NVIDIA产品的质量和可靠性，并可能导致超出本文档范围的其他或不同的条件和/或要求。对于任何因以下原因导致的故障、损坏、成本或问题，NVIDIA不承担任何责任：(i) 以违反本文档的任何方式使用NVIDIA产品或(ii) 客户产品设计。
本文档不授予任何NVIDIA专利权、版权或其他NVIDIA知识产权的明示或暗示许可。NVIDIA发布的关于第三方产品或服务的信息，不构成NVIDIA对这些产品或服务的使用许可或担保认可。使用此类信息可能需要获得第三方基于其专利或其他知识产权的许可，或需要获得NVIDIA基于其专利或其他知识产权的许可。
本文件中的信息仅可在获得NVIDIA事先书面批准、未经改动完整复制且完全符合所有适用的出口法律法规，并附带所有相关条件、限制和声明的情况下进行复制。
本文件及所有NVIDIA设计规格、参考板、文件、图纸、诊断工具、清单和其他文档（统称及单独称为"材料"）均以"现状"提供。NVIDIA不对材料作出任何明示或默示的保证，包括但不限于对不侵权、适销性和特定用途适用性的默示保证免责。在法律允许的最大范围内，NVIDIA不就因使用本文件导致的任何损害承担责任，包括但不限于任何直接、间接、特殊、附带、惩罚性或后果性损害，无论损害成因如何，也无论责任理论为何，即使NVIDIA已被告知发生此类损害的可能性。不论客户因任何原因可能遭受的任何损害，NVIDIA对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。
8.2. OpenCL
OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。
8.3. 商标
NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。