NVIDIA CUDA C ++编译器的新特性
cuda 11 . 5 c ++编译器解决了不断增长的客户请求。具体来说，如何减少 cuda 应用程序构建时间。除了消除未使用的内核外， nvrtc 和 ptx 并发编译有助于解决这个关键问题 cuda c ++应用程序开发的关注点。
cuda 11 . 5 nvcc 编译器现在添加了对 clang 12 . 0 作为主机编译器的支持。我们还提供了 128 位整数支持的有限预览版本，这在高保真计算中变得至关重要。
cuda c ++编译器工具链上的技术演练补充了编程指南（需要链接），并提供了在 cuda 11 . 5 工具包版本中引入的新特性的广泛概述。
并发编译 nvrtc 编译过程分为三个主要阶段：
parser -> nvvm optimizer -> ptx compiler 其中一些阶段不是线程安全的，因此 nvrtc 以前会使用全局锁序列化来自多个用户线程的并发编译请求。
在 cuda 11 . 5 中，对 nvrtc 实现进行了增强，以提供部分并发编译支持。这是通过移除全局锁和使用每阶段锁来实现的，这会导致不同的线程并发执行编译管道的不同阶段。
图 1 显示了 cuda 11 . 5 之前的 nvrtc 如何序列化来自四个线程的同时编译请求。
图 1 。序列化编译
对于 11 . 5 ， nvrtc 不会序列化编译请求。相反，来自不同线程的编译请求是管道化的，从而使编译管道的不同阶段能够同时进行。
图 2 。并发编译
图 3 中的图表显示了编译一组 100 个相同的示例 nvrtc 程序的总编译时间，这些程序按可用线程数进行划分。
图 3 。 cuda 11 . 4 和 11 . 5 的编译时间比较
正如所料，对于 cuda 11 . 4 nvrtc ，总编译时间不会随着线程数的增加而改变，而编译是使用全局 nvrtc 锁序列化的。使用 cuda 11 . 5 nvrtc ，总编译时间会随着线程数的增加而减少。我们将继续使单个 stage 线程更安全，这将为本例实现近乎线性的加速。
ptx 并发编译 沿着 jit 编译路径进行 ptx 编译，以及使用 ptx 静态库进行多个内部阶段。这些阶段以前的实现不能保证从多个线程进行并发编译。相反， ptx 编译器使用全局锁来序列化并发编译。
在 cuda 11 . 5 和 r495 驱动程序中， ptx 编译器实现现在使用更细粒度的本地锁，而不是全局锁。这允许并发执行多个编译请求，并显著缩短了编译时间。
下图显示了编译 104 个相同的示例程序所需的总编译时间，这些程序在给定数量的线程上拆分到culinkadddata使用cu_jit_input_ptx作为cujitinputtype。
正如 r470 cuda 驱动程序所预期的那样，总编译时间不会随着线程数的增加而改变，因为编译是用全局锁序列化的。使用 r495 cuda 驱动程序，总编译时间随着线程数的增加而减少。
图 4 。 cuda 11 . 4 和 11 . 5 的 ptx 并发编译比较
消除未使用的内核 单独编译模式允许 cuda 内核函数和设备函数作为 cuda 设备代码库发布，并使用设备链接器 nvlink 针对任何用户应用程序进行链接。然后在运行时在 gpu 上加载并执行生成的设备程序。
在 cuda 11 . 5 之前， nvlink 无法确定从链接设备程序中删除未使用的内核是否安全，因为这些内核函数可以从主机代码中引用。
考虑一个定义四个内核函数的库：
//library.cu__global__ void aaa() { /* code */ }__global__ void bbb() { /* code */ }__global__ void ccc() { /* code */ }__global__ void ddd() { /* code */ }
该库的构建和发布：
$nvcc -rdc=true library.cu -lib -o testlib.a
用户代码引用库中的单个内核：
//user.cuextern __global__ void aaa();int main() { aaa<>(); }
代码链接为：
$nvcc -rdc=true user.cu testlib.a -o user
以 cuda 11 . 4 为例，链接设备程序将包含所有四个内核体，即使链接设备程序中只使用一个内核（“ aaa ”）。对于链接到较大库的应用程序来说，这可能是一个负担。
增加的二进制大小和应用程序加载时间并不是冗余设备代码的唯一问题。当使用设备链接时间优化( dlto –修复链接)时，在优化之前未删除的未使用内核可能会导致更长的构建时间，并可能阻碍代码优化。
使用 cuda 11 . 5 ， cuda 编译器将跟踪主机代码中对内核的引用，并将此信息传播到设备链接器（ nvlink ）。 nvlink 然后从链接的设备程序中删除未使用的内核。对于前面的示例，未使用的内核 bbb 、 ccc 和 ddd 将从链接设备程序中删除。
在 cuda 11 . 5 中，默认情况下禁用此优化，但可以通过将-xnvlink -use-host-info选项添加到 nvcc 命令行来启用：
$nvcc -rdc=true user.cu testlib.a -o user -xnvlink -use-host-info
在随后的 cuda 工具包版本中，默认情况下将启用优化，并提供一个退出标志。
这里有一些警告。在 cuda 11 . 5 中，编译器对内核引用的分析在以下情况下是保守的。编译器可以考虑一些未从宿主代码实际引用的内核，如：
如果模板实例化是从主机代码引用的，则该模板的所有实例都被视为是从主机代码引用的。template__global__ void foo() { }__device__ void doit() { foo<>(); } int main() {// compiler will mark all instances of foo template as referenced// from host code, including foo, which is only actually // referenced from device codefoo<>();} __global__ or __device__函数体之外的任何引用都被视为主机代码引用。__global__ void foo() { }__device__ auto *ptr = foo; // foo is considered as referenced // from host code. 当对函数的引用为template-dependent时，具有该名称的所有内核都被视为主机引用。__global__ void foo(int) { }namespace n1 {template __global__ void foo(t) { }}templatevoid doit() { // the reference to 'foo' is template dependent, so // both ::foo and all instances of ::n1::foo are // considered as referenced from host code. foo<>(t{});}
另一个警告是，当设备链接步骤推迟到主机应用程序启动（ jit 链接）时，而不是在构建时，将不会删除未使用的内核。
// with nonvirtual architecture (sm_80), nvlink is invoked // at build time, and kernel pruning will occur.$nvcc -xnvlink -use-host-info -rdc=true foo.cu bar.cu -o foo -arch sm_80// with virtual architecture (compute_80), nvlink is not invoked// at build time, but only during host application startup.// kernel pruning will not occur.$nvcc -xnvlink -use-host-info -rdc=true foo.cu bar.cu -o foo -arch compute_80
今后的工作 在 cuda 11 . 5 中， nvlink 在设备链接时间优化（ dlto – fixme link ）期间尚未使用有关未使用内核的信息。我们的目标是使 nvlink 能够使用此信息删除未使用的内核，减少优化器时间，并通过减少代码膨胀来提高生成的代码质量。
有限的 128 位整数支持 11 . 5 cuda c ++编译器支持主机编译器支持 128 位整数的平台的 128 位整数数据类型。基本的算术、逻辑和位运算将在 128 位整数上工作。未来版本计划支持 cuda 固有类型和 cuda 数学函数的 128 位整数变体。
类似地，对 128 位整数的调试支持以及与开发人员工具的集成将在后续版本中提供。目前，我们正在开发者论坛上寻求您对此预览功能的早期反馈。
nvrtc 静态库 cuda 11 . 5 提供了 nvrtc 库的静态版本。一些应用程序可能更喜欢链接静态 nvrtc 库，以保证部署期间的稳定性能和功能。静态库用户还希望静态链接 nvrtc 内置库和 ptx 编译器库的静态版本。有关链接静态 nvrtc 库的更多信息，请参阅nvrtc 用户指南。
__builtin_assume cuda 11 . 5 改进了__builtin_assume应用于__isshared(pointer)等地址空间谓词函数的结果时加载和存储的代码生成。有关其他支持的功能，请参阅地址空间谓词函数。
如果没有地址空间说明符，编译器将生成通用加载和存储指令，这需要一些额外的指令来计算特定的内存段，然后再执行实际的内存操作。使用__builtin_assume(expr)提示编译器使用泛型指针的地址空间，这可能会提高代码的性能。
correct usage: bool b = __isshared(ptr); __builtin_assume(b); // ok: proof that ptr is a pointer to shared memoryincorrect usage:these hints are ignored unless the boolean expression is stored in a separate variable: __builtin_assume(__isshared(ptr)); // ignored
与其他__builtin_assume一样，如果表达式不为 true ，则行为未定义。如果您有兴趣了解__builtin_assume的更多信息，请参阅 cuda 11 . 2 编译器文章。
pragma 诊断控制 在 cuda 11 . 5 中， nvcc cuda 编译器前端增加了对大量杂注的支持，这些杂注提供了对诊断消息的更多控制。
您可以使用以下杂注来控制特定错误号的编译器诊断：
#pragma nv_diag_suppress // suppress the specified diagnostic // message#pragma nv_diag_warning // make the specified diagnostic a warning#pragma nv_diag_error // make the specified diagnostic an error#pragma nv_diag_default // restore the specified diagnostic level // to default#pragma nv_diag_once // only report the specified diagnostic onceuses of these pragmas have the following form:#pragma nv_diag_xxx error_number, error_number …
要了解如何使用这些带有更详细警告的杂注，请参阅 cuda 内部编程指南。以下示例将取消foo的 clara 选项上的“declared but never referenced”警告：
#pragma nv_diag_suppress 177void foo(){ int xxx=0;}
杂注nv_diagnostic推送和nv_diagnostic弹出可用于保存和恢复当前诊断pragma状态：
#pragma nv_diagnostic push#pragma nv_diag_suppress 177void foo(){ int xxx=0;}#pragma nv_diagnostic popvoid bar(){ int xxx=0;}
这些杂注都不会对主机编译器产生任何影响。
不推荐使用注意：不带nv_前缀的诊断杂注已不推荐使用。例如，#pragma diag_suppress支持将从所有未来版本中删除。使用这些诊断标记将引发如下警告消息：
pragma diag_suppress is deprecated, use nv_diag_suppress instead
宏__nvcc_diag_pragma_support__有助于过渡到使用新宏：
#ifdef __nvcc_diag_pragma_support__#pragma nv_diag_suppress 177#else#pragma diag_suppress 177#endif
新选项 -arch = all | all major 在 cuda 11 . 5 版本之前，如果您想为所有受支持的体系结构生成代码，必须在--generate-code选项中列出所有目标。如果添加了较新的版本，或旧版本失效，则必须相应地更改--generate-code选项。现在，新选项-arch=all|all-major提供了一种更简单、更高效的方法。
如果指定了-arch=all， nvcc 将为所有受支持的体系结构(sm_*)嵌入编译后的代码映像，并为最高的主要虚拟体系结构嵌入 ptx 程序。
如果指定了-arch=all-major， nvcc 将为所有受支持的主要版本(sm_*0)嵌入编译后的代码映像，从最早受支持的sm_x体系结构（此版本为sm_35）开始，并为最高的主要虚拟体系结构嵌入 ptx 程序。
例如，简单的-arch=all选项相当于此版本的以下一长串选项：
-gencode arch=compute_35,\code=sm_35\ -gencode arch=compute_37,\code=sm_37\ -gencode arch=compute_50,\code=sm_50\ -gencode arch=compute_52,\code=sm_52\ -gencode arch=compute_53,\code=sm_53\-gencode arch=compute_60,\code=sm_60\ -gencode arch=compute_61,\code=sm_61\ -gencode arch=compute_62,\code=sm_62\ -gencode arch=compute_70,\code=sm_70\ -gencode arch=compute_72,\code=sm_72\ -gencode arch=compute_75,\code=sm_75\ -gencode arch=compute_80,\code=sm_80\ -gencode arch=compute_86,\code=sm_86\ -gencode arch=compute_87,\code=sm_87\ -gencode arch=compute_80,\code=compute_80\
简单的-arch=all-major选项相当于此版本的以下一长串选项：
-gencode arch=compute_35,\code=sm_35\ -gencode arch=compute_50,\code=sm_50\ -gencode arch=compute_60,\code=sm_60\ -gencode arch=compute_70,\code=sm_70\ -gencode arch=compute_80,\code=sm_80\ -gencode arch=compute_80,\code=compute_80\
有关所有受支持的虚拟体系结构，请参阅虚拟体系结构功能列表。有关所有受支持的真实体系结构，请参阅 gpu 功能列表。
确定性代码生成 在以前的 cuda 工具包中，设备代码中内部链接变量或函数的名称在每次 nvcc 调用时都会更改，即使源代码没有更改。某些软件管理和构建系统检查生成的程序位是否已更改。先前的 nvcc 编译器行为导致此类系统触发，并错误地假设源程序中存在语义更改；例如，可能触发冗余的依赖生成。
在 cuda 11 . 5 中， nvcc 编译器行为已更改为确定性。例如，考虑这个测试用例：
//--static __device__ void foo() { }auto __device__ fptr = foo;int main() { }//--
在 cuda 11 . 4 中，两次编译同一程序会在 ptx 中生成稍微不同的名称：
//--$cuda-11.4/bin/nvcc -std=c++14 -rdc=true -ptx test.cu -o test1.ptx$cuda-11.4/bin/nvcc -std=c++14 -rdc=true -ptx test.cu -o test2.ptx$diff -w test1.ptx test2.ptx13c13 .func _zn57_internal_39_tmpxft_00000a4e_00000000_7_test_cpp1_ii_main3fooev16c16 .visible .global .align 8 .u64 fptr = _zn57_internal_39_tmpxft_00000a4e_00000000_7_test_cpp1_ii_main3fooev;18c18 .func _zn57_internal_39_tmpxft_00000a4e_00000000_7_test_cpp1_ii_main3fooev()$//--
使用 cuda 11 . 5 ，两次编译同一程序会生成相同的 ptx ：
//--$nvcc -std=c++14 -rdc=true -ptx test.cu -o test1.ptx$nvcc -std=c++14 -rdc=true -ptx test.cu -o test2.ptx $diff -w test1.ptx test2.ptx$//--
结论
通过阅读在 cuda 11 . 5 工具包中展示新功能文章，了解更多关于 cuda 11 . 5 工具包的信息。
关于作者
arthy sundaram 是 cuda 平台的技术产品经理。她拥有哥伦比亚大学计算机科学硕士学位。她感兴趣的领域是操作系统、编译器和计算机体系结构。
jaydeep marathe 是 nvidia 的高级编译工程师。他拥有北卡罗来纳州立大学计算机科学硕士和博士学位。
hari sandanagobalane 是 nvidia 的高级编译工程师。他拥有新加坡国立大学计算机科学硕士学位。
mike murphy 是 nvidia 的高级编译工程师。
xiaohua zhang 是 nvidia 的高级编译工程师。他拥有清华大学计算机科学硕士学位。
girish bharambe 是 nvidia 的高级编译经理。他拥有印度浦那大学计算机工程学士学位。
2023年财富中国500强：国家电网居首 比亚迪营收增长超九成
Starlink卫星宽带未来将提供1Gbps每用户的网速
光I/O接口革命：4Tb/s超高速传输引领未来
美国JDAM联合制导攻击武器技术解析与实战应用
小米路由器AC2100高清图集
NVIDIA CUDA C ++编译器的新特性
国产化显示器接口芯片：Type-C接口台式显示器方案
京东空调全渠道发力深挖市场增长 助力行业持续深化变革
具有高可靠性的单膜电容式微机械超声换能器（CMUT）阵列元件
如果项目的模型遇到瓶颈，用这些Tricks就对了
雷军在微博称，小米9拥有3999元以内手机“最美下巴”
一文详解磁敏传感器技术原理
阿朗创新型数据中心交换解决方案荣获Best of Interop大奖
Win10如何助力微软一统物联网生态？
2020年第三季度全球网络硬件收入同比Q1下降3%
ANYMESH-SDR-A4室外固定基站型自组网设备介绍
2018年国内电动汽车十大投资回顾 特斯拉成第一家外商独资车企
哪种液位传感器不易产生污垢
LED灯关灯后闪烁的解决方法
苹果祸不单行 iPhone 15没往年那么难抢 iPhone 12辐射超标?多国跟进调查
cuda 11 . 5 nvcc 编译器现在添加了对 clang 12 . 0 作为主机编译器的支持。我们还提供了 128 位整数支持的有限预览版本，这在高保真计算中变得至关重要。
cuda c ++编译器工具链上的技术演练补充了编程指南（需要链接），并提供了在 cuda 11 . 5 工具包版本中引入的新特性的广泛概述。
并发编译 nvrtc 编译过程分为三个主要阶段：
parser -> nvvm optimizer -> ptx compiler 其中一些阶段不是线程安全的，因此 nvrtc 以前会使用全局锁序列化来自多个用户线程的并发编译请求。
在 cuda 11 . 5 中，对 nvrtc 实现进行了增强，以提供部分并发编译支持。这是通过移除全局锁和使用每阶段锁来实现的，这会导致不同的线程并发执行编译管道的不同阶段。
图 1 显示了 cuda 11 . 5 之前的 nvrtc 如何序列化来自四个线程的同时编译请求。
图 1 。序列化编译
对于 11 . 5 ， nvrtc 不会序列化编译请求。相反，来自不同线程的编译请求是管道化的，从而使编译管道的不同阶段能够同时进行。
图 2 。并发编译
图 3 中的图表显示了编译一组 100 个相同的示例 nvrtc 程序的总编译时间，这些程序按可用线程数进行划分。
图 3 。 cuda 11 . 4 和 11 . 5 的编译时间比较
正如所料，对于 cuda 11 . 4 nvrtc ，总编译时间不会随着线程数的增加而改变，而编译是使用全局 nvrtc 锁序列化的。使用 cuda 11 . 5 nvrtc ，总编译时间会随着线程数的增加而减少。我们将继续使单个 stage 线程更安全，这将为本例实现近乎线性的加速。
ptx 并发编译 沿着 jit 编译路径进行 ptx 编译，以及使用 ptx 静态库进行多个内部阶段。这些阶段以前的实现不能保证从多个线程进行并发编译。相反， ptx 编译器使用全局锁来序列化并发编译。
在 cuda 11 . 5 和 r495 驱动程序中， ptx 编译器实现现在使用更细粒度的本地锁，而不是全局锁。这允许并发执行多个编译请求，并显著缩短了编译时间。
下图显示了编译 104 个相同的示例程序所需的总编译时间，这些程序在给定数量的线程上拆分到culinkadddata使用cu_jit_input_ptx作为cujitinputtype。
正如 r470 cuda 驱动程序所预期的那样，总编译时间不会随着线程数的增加而改变，因为编译是用全局锁序列化的。使用 r495 cuda 驱动程序，总编译时间随着线程数的增加而减少。
图 4 。 cuda 11 . 4 和 11 . 5 的 ptx 并发编译比较
消除未使用的内核 单独编译模式允许 cuda 内核函数和设备函数作为 cuda 设备代码库发布，并使用设备链接器 nvlink 针对任何用户应用程序进行链接。然后在运行时在 gpu 上加载并执行生成的设备程序。
在 cuda 11 . 5 之前， nvlink 无法确定从链接设备程序中删除未使用的内核是否安全，因为这些内核函数可以从主机代码中引用。
考虑一个定义四个内核函数的库：
//library.cu__global__ void aaa() { /* code */ }__global__ void bbb() { /* code */ }__global__ void ccc() { /* code */ }__global__ void ddd() { /* code */ }
该库的构建和发布：
$nvcc -rdc=true library.cu -lib -o testlib.a
用户代码引用库中的单个内核：
//user.cuextern __global__ void aaa();int main() { aaa<>(); }
代码链接为：
$nvcc -rdc=true user.cu testlib.a -o user
以 cuda 11 . 4 为例，链接设备程序将包含所有四个内核体，即使链接设备程序中只使用一个内核（“ aaa ”）。对于链接到较大库的应用程序来说，这可能是一个负担。
增加的二进制大小和应用程序加载时间并不是冗余设备代码的唯一问题。当使用设备链接时间优化( dlto –修复链接)时，在优化之前未删除的未使用内核可能会导致更长的构建时间，并可能阻碍代码优化。
使用 cuda 11 . 5 ， cuda 编译器将跟踪主机代码中对内核的引用，并将此信息传播到设备链接器（ nvlink ）。 nvlink 然后从链接的设备程序中删除未使用的内核。对于前面的示例，未使用的内核 bbb 、 ccc 和 ddd 将从链接设备程序中删除。
在 cuda 11 . 5 中，默认情况下禁用此优化，但可以通过将-xnvlink -use-host-info选项添加到 nvcc 命令行来启用：
$nvcc -rdc=true user.cu testlib.a -o user -xnvlink -use-host-info
在随后的 cuda 工具包版本中，默认情况下将启用优化，并提供一个退出标志。
这里有一些警告。在 cuda 11 . 5 中，编译器对内核引用的分析在以下情况下是保守的。编译器可以考虑一些未从宿主代码实际引用的内核，如：
如果模板实例化是从主机代码引用的，则该模板的所有实例都被视为是从主机代码引用的。template__global__ void foo() { }__device__ void doit() { foo<>(); } int main() {// compiler will mark all instances of foo template as referenced// from host code, including foo, which is only actually // referenced from device codefoo<>();} __global__ or __device__函数体之外的任何引用都被视为主机代码引用。__global__ void foo() { }__device__ auto *ptr = foo; // foo is considered as referenced // from host code. 当对函数的引用为template-dependent时，具有该名称的所有内核都被视为主机引用。__global__ void foo(int) { }namespace n1 {template __global__ void foo(t) { }}templatevoid doit() { // the reference to 'foo' is template dependent, so // both ::foo and all instances of ::n1::foo are // considered as referenced from host code. foo<>(t{});}
另一个警告是，当设备链接步骤推迟到主机应用程序启动（ jit 链接）时，而不是在构建时，将不会删除未使用的内核。
// with nonvirtual architecture (sm_80), nvlink is invoked // at build time, and kernel pruning will occur.$nvcc -xnvlink -use-host-info -rdc=true foo.cu bar.cu -o foo -arch sm_80// with virtual architecture (compute_80), nvlink is not invoked// at build time, but only during host application startup.// kernel pruning will not occur.$nvcc -xnvlink -use-host-info -rdc=true foo.cu bar.cu -o foo -arch compute_80
今后的工作 在 cuda 11 . 5 中， nvlink 在设备链接时间优化（ dlto – fixme link ）期间尚未使用有关未使用内核的信息。我们的目标是使 nvlink 能够使用此信息删除未使用的内核，减少优化器时间，并通过减少代码膨胀来提高生成的代码质量。
有限的 128 位整数支持 11 . 5 cuda c ++编译器支持主机编译器支持 128 位整数的平台的 128 位整数数据类型。基本的算术、逻辑和位运算将在 128 位整数上工作。未来版本计划支持 cuda 固有类型和 cuda 数学函数的 128 位整数变体。
类似地，对 128 位整数的调试支持以及与开发人员工具的集成将在后续版本中提供。目前，我们正在开发者论坛上寻求您对此预览功能的早期反馈。
nvrtc 静态库 cuda 11 . 5 提供了 nvrtc 库的静态版本。一些应用程序可能更喜欢链接静态 nvrtc 库，以保证部署期间的稳定性能和功能。静态库用户还希望静态链接 nvrtc 内置库和 ptx 编译器库的静态版本。有关链接静态 nvrtc 库的更多信息，请参阅nvrtc 用户指南。
__builtin_assume cuda 11 . 5 改进了__builtin_assume应用于__isshared(pointer)等地址空间谓词函数的结果时加载和存储的代码生成。有关其他支持的功能，请参阅地址空间谓词函数。
如果没有地址空间说明符，编译器将生成通用加载和存储指令，这需要一些额外的指令来计算特定的内存段，然后再执行实际的内存操作。使用__builtin_assume(expr)提示编译器使用泛型指针的地址空间，这可能会提高代码的性能。
correct usage: bool b = __isshared(ptr); __builtin_assume(b); // ok: proof that ptr is a pointer to shared memoryincorrect usage:these hints are ignored unless the boolean expression is stored in a separate variable: __builtin_assume(__isshared(ptr)); // ignored
与其他__builtin_assume一样，如果表达式不为 true ，则行为未定义。如果您有兴趣了解__builtin_assume的更多信息，请参阅 cuda 11 . 2 编译器文章。
pragma 诊断控制 在 cuda 11 . 5 中， nvcc cuda 编译器前端增加了对大量杂注的支持，这些杂注提供了对诊断消息的更多控制。
您可以使用以下杂注来控制特定错误号的编译器诊断：
#pragma nv_diag_suppress // suppress the specified diagnostic // message#pragma nv_diag_warning // make the specified diagnostic a warning#pragma nv_diag_error // make the specified diagnostic an error#pragma nv_diag_default // restore the specified diagnostic level // to default#pragma nv_diag_once // only report the specified diagnostic onceuses of these pragmas have the following form:#pragma nv_diag_xxx error_number, error_number …
要了解如何使用这些带有更详细警告的杂注，请参阅 cuda 内部编程指南。以下示例将取消foo的 clara 选项上的“declared but never referenced”警告：
#pragma nv_diag_suppress 177void foo(){ int xxx=0;}
杂注nv_diagnostic推送和nv_diagnostic弹出可用于保存和恢复当前诊断pragma状态：
#pragma nv_diagnostic push#pragma nv_diag_suppress 177void foo(){ int xxx=0;}#pragma nv_diagnostic popvoid bar(){ int xxx=0;}
这些杂注都不会对主机编译器产生任何影响。
不推荐使用注意：不带nv_前缀的诊断杂注已不推荐使用。例如，#pragma diag_suppress支持将从所有未来版本中删除。使用这些诊断标记将引发如下警告消息：
pragma diag_suppress is deprecated, use nv_diag_suppress instead
宏__nvcc_diag_pragma_support__有助于过渡到使用新宏：
#ifdef __nvcc_diag_pragma_support__#pragma nv_diag_suppress 177#else#pragma diag_suppress 177#endif
新选项 -arch = all | all major 在 cuda 11 . 5 版本之前，如果您想为所有受支持的体系结构生成代码，必须在--generate-code选项中列出所有目标。如果添加了较新的版本，或旧版本失效，则必须相应地更改--generate-code选项。现在，新选项-arch=all|all-major提供了一种更简单、更高效的方法。
如果指定了-arch=all， nvcc 将为所有受支持的体系结构(sm_*)嵌入编译后的代码映像，并为最高的主要虚拟体系结构嵌入 ptx 程序。
如果指定了-arch=all-major， nvcc 将为所有受支持的主要版本(sm_*0)嵌入编译后的代码映像，从最早受支持的sm_x体系结构（此版本为sm_35）开始，并为最高的主要虚拟体系结构嵌入 ptx 程序。
例如，简单的-arch=all选项相当于此版本的以下一长串选项：
-gencode arch=compute_35,\code=sm_35\ -gencode arch=compute_37,\code=sm_37\ -gencode arch=compute_50,\code=sm_50\ -gencode arch=compute_52,\code=sm_52\ -gencode arch=compute_53,\code=sm_53\-gencode arch=compute_60,\code=sm_60\ -gencode arch=compute_61,\code=sm_61\ -gencode arch=compute_62,\code=sm_62\ -gencode arch=compute_70,\code=sm_70\ -gencode arch=compute_72,\code=sm_72\ -gencode arch=compute_75,\code=sm_75\ -gencode arch=compute_80,\code=sm_80\ -gencode arch=compute_86,\code=sm_86\ -gencode arch=compute_87,\code=sm_87\ -gencode arch=compute_80,\code=compute_80\
简单的-arch=all-major选项相当于此版本的以下一长串选项：
-gencode arch=compute_35,\code=sm_35\ -gencode arch=compute_50,\code=sm_50\ -gencode arch=compute_60,\code=sm_60\ -gencode arch=compute_70,\code=sm_70\ -gencode arch=compute_80,\code=sm_80\ -gencode arch=compute_80,\code=compute_80\
有关所有受支持的虚拟体系结构，请参阅虚拟体系结构功能列表。有关所有受支持的真实体系结构，请参阅 gpu 功能列表。
确定性代码生成 在以前的 cuda 工具包中，设备代码中内部链接变量或函数的名称在每次 nvcc 调用时都会更改，即使源代码没有更改。某些软件管理和构建系统检查生成的程序位是否已更改。先前的 nvcc 编译器行为导致此类系统触发，并错误地假设源程序中存在语义更改；例如，可能触发冗余的依赖生成。
在 cuda 11 . 5 中， nvcc 编译器行为已更改为确定性。例如，考虑这个测试用例：
//--static __device__ void foo() { }auto __device__ fptr = foo;int main() { }//--
在 cuda 11 . 4 中，两次编译同一程序会在 ptx 中生成稍微不同的名称：
//--$cuda-11.4/bin/nvcc -std=c++14 -rdc=true -ptx test.cu -o test1.ptx$cuda-11.4/bin/nvcc -std=c++14 -rdc=true -ptx test.cu -o test2.ptx$diff -w test1.ptx test2.ptx13c13 .func _zn57_internal_39_tmpxft_00000a4e_00000000_7_test_cpp1_ii_main3fooev16c16 .visible .global .align 8 .u64 fptr = _zn57_internal_39_tmpxft_00000a4e_00000000_7_test_cpp1_ii_main3fooev;18c18 .func _zn57_internal_39_tmpxft_00000a4e_00000000_7_test_cpp1_ii_main3fooev()$//--
使用 cuda 11 . 5 ，两次编译同一程序会生成相同的 ptx ：
//--$nvcc -std=c++14 -rdc=true -ptx test.cu -o test1.ptx$nvcc -std=c++14 -rdc=true -ptx test.cu -o test2.ptx $diff -w test1.ptx test2.ptx$//--
结论
通过阅读在 cuda 11 . 5 工具包中展示新功能文章，了解更多关于 cuda 11 . 5 工具包的信息。
关于作者
arthy sundaram 是 cuda 平台的技术产品经理。她拥有哥伦比亚大学计算机科学硕士学位。她感兴趣的领域是操作系统、编译器和计算机体系结构。
jaydeep marathe 是 nvidia 的高级编译工程师。他拥有北卡罗来纳州立大学计算机科学硕士和博士学位。
hari sandanagobalane 是 nvidia 的高级编译工程师。他拥有新加坡国立大学计算机科学硕士学位。
mike murphy 是 nvidia 的高级编译工程师。
xiaohua zhang 是 nvidia 的高级编译工程师。他拥有清华大学计算机科学硕士学位。
girish bharambe 是 nvidia 的高级编译经理。他拥有印度浦那大学计算机工程学士学位。
2023年财富中国500强：国家电网居首 比亚迪营收增长超九成
Starlink卫星宽带未来将提供1Gbps每用户的网速
光I/O接口革命：4Tb/s超高速传输引领未来
美国JDAM联合制导攻击武器技术解析与实战应用
小米路由器AC2100高清图集
NVIDIA CUDA C ++编译器的新特性
国产化显示器接口芯片：Type-C接口台式显示器方案
京东空调全渠道发力深挖市场增长 助力行业持续深化变革
具有高可靠性的单膜电容式微机械超声换能器（CMUT）阵列元件
如果项目的模型遇到瓶颈，用这些Tricks就对了
雷军在微博称，小米9拥有3999元以内手机“最美下巴”
一文详解磁敏传感器技术原理
阿朗创新型数据中心交换解决方案荣获Best of Interop大奖
Win10如何助力微软一统物联网生态？
2020年第三季度全球网络硬件收入同比Q1下降3%
ANYMESH-SDR-A4室外固定基站型自组网设备介绍
2018年国内电动汽车十大投资回顾 特斯拉成第一家外商独资车企
哪种液位传感器不易产生污垢
LED灯关灯后闪烁的解决方法
苹果祸不单行 iPhone 15没往年那么难抢 iPhone 12辐射超标?多国跟进调查