CUDA C++ 编译器更新对 ELF 可见性和链接影响

在下一个 CUDA 主要版本 CUDA 13.0 中，NVIDIA 将对 NVIDIA CUDA 编译器驱动程序 (NVCC) 作出两项重大更改，这些更改将影响 __global__ 函数和设备变量的 ELF 可见性和关联。这些更新旨在防止长期以来难以检测和调试的细微运行时错误。但是，这些更改可能会影响一些现有的 CUDA C++ 程序。

本文旨在提醒用户注意潜在的中断，解释更改背后的原因，并就可以恢复旧版行为的 NVCC 标志提供指导。表 1 汇总了这两项更改。

NVCC 变化# 1：ELF 可见性

在 CUDA 13.0 之前的工具包中，NVCC 编译器未修改发送给主机编译器的代码中 __global__ 函数和 __managed__/__device__/__constant__ 变量的 ELF 可见性。如果将生成的代码打包到共享库中，这些符号将对共享库的用户可见。

带示例的问题概述

默认情况下，NVCC 链接 CUDA Runtime Library (CUDART) 的静态版本。这将导致两个不同的 CUDART 库链接：一个连接到共享库，另一个连接到主程序。如果通过共享库边界访问 __global__ 核函数或 __device__/__managed__/__constant__ 变量，则可能会导致细微的运行时问题 (图 1) 。

示例 1

//-- foo.cu --
#include <cstdio>
__global__ void foo() {   
printf("\n hi!"); 
}

//-- main.cu --
#include <cstdio>
extern __global__ void foo();
int main() {
  foo<<<1,1>>>();
  cudaDeviceSynchronize();
  auto err = cudaGetLastError();
  printf("\n cudaGetLastError() = %s\n", cudaGetErrorString(err));
}

foo.cu 内置于共享库 libfoo.so 中，并从主程序中引用：

$nvcc foo.cu -shared -o libfoo.so -Xcompiler -fPIC -rdc=true
$nvcc main.cu libfoo.so -o main -rdc=true

运行此程序时，预期行 (“hi”) 未打印出来，但 CUDA Runtime 未报告任何错误：

LD_LIBRARY_PATH=. ./main                                                                   
 cudaGetLastError() = no error

潜在问题是，__global__ 函数核函数启动序列涉及在 main.cu (foo<<<...>>>) 中的启动点以及 foo.cu 中 foo 的主机代码存根函数内部调用 CUDA Runtime (例如，用于打包任何函数参数) 。但是，由于 libfoo.so 和 main 程序中链接了不同的 CUDART 库，因此内核启动无法达到预期效果。

示例 2

//foo.cu
__managed__ int result = 20;

//main.cu
#include <cstdio>
extern __managed__ int result;
int main() {
  printf("\n result = %d", result);
}

与第一个示例一样，foo.cu 内置于共享库 libfoo.so 中，并引用自主程序：

$nvcc foo.cu -shared -o libfoo.so -Xcompiler -fPIC -rdc=true
$nvcc main.cu libfoo.so -o main -rdc=true -g

程序运行时，在访问 result 的值时，会在 main.cu 中遇到分割错误。

$LD_LIBRARY_PATH=. gdb -ex=r ./main
..
Thread 1 "main" received signal SIGSEGV, Segmentation fault.
0x000055555555cdaf in main () at main.cu:4
4         printf("\n result = %d", result);

同样，潜在的问题是不同的 CUDART 库以静态方式连接到 libfoo.so 和 main 程序，这会干扰 __managed__ 变量 result 的正确初始化。

这些示例展示了从共享库中导出 __global__ 函数或 __managed__ 变量符号时的运行时崩溃和细微的意外运行时行为。这些问题难以追踪 (没有构建时警告或运行时 CUDA 错误) ，并且涉及一个或多个共享库 (可能由不同供应商提供) 与主程序之间的交互。

受影响的平台

本节所述的 CUDA 13.0 NVCC 更改会影响除 Windows 以外的所有平台。默认情况下，Windows 上的主机编译器工具链 (cl.exe) 不会从共享库中导出符号，因此不会出现本节所述的问题。

CUDA 13.0 中引入的解决方案

为避免用户遇到上述问题，从 CUDA 13.0 开始，NVCC 将 __global__ 函数和 __managed__/__device__/__constant__ 变量的默认可见性更改为 hidden，导致此类符号在共享库之外不再可见。

这将导致上述程序构建失败，这比运行时崩溃或错误行为更可取。例如，对于上面的 __managed__ 变量示例，将出现链接失败：

$nvcc foo.cu -shared -o libfoo.so -Xcompiler -fPIC -rdc=true 
$nvcc main.cu libfoo.so -o main -rdc=true   -g 
...
/usr/bin/ld: /tmp/tmpxft_0032ad7a_00000000-11_main.o: in function `main':
/work/bugs/blogexamples/sharedlibrary/managed_var/main.cu:4: undefined reference to `result'

在构建 主程序和共享库 时，通过在 CUDART 的共享库版本 (-cudart=shared) 中持续链接，可以避免正确性问题。通过这种方法，程序的所有部分都使用相同的 CUDART。然后禁用 NVCC 强制隐藏可见性 (--device-entity-has-hidden-visibility=false)，构建如下：

$nvcc foo.cu -shared -o libfoo.so -Xcompiler -fPIC -rdc=true -cudart=shared -device-entity-has-hidden-visibility=false
$nvcc main.cu libfoo.so -o main -rdc=true   -g -cudart=shared -device-entity-has-hidden-visibility=false

选择退出 ( CUDA 13.0)

我们认识到，默认 NVCC 行为的这种变化可能会破坏一些现有工作流程。特别是，它可能会影响工作流程，因为共享库和使用 -cudart=shared NVCC 标志的主程序仅使用一个动态 (共享) CUDART 库。

要在 CUDA 13.0 及更高版本中继续支持这些工作流程，您可以选择退出已更改的 NVCC 行为。选择退出：

• 将 --device-entity-has-hidden-visibility=false 添加到 NVCC 命令行中。 此操作可将 NVCC 行为恢复到前 CUDA 13.0 工具包的行为。请注意，此标志自 CUDA 12.8 起可用，但在 CUDA 13.0 之前，默认值为 false。
• 在 __global__ 函数或 __managed__/__device__/__constant__ 变量声明、封闭命名空间或使用 #pragma GCC visibility 显式添加 __attribute__((visibility("default")))。

有关详细信息，请参阅以下示例：

__global__ __attribute__((visibility("default"))) void foo1() { }
namespace __attribute__((visibility("default"))) N1 { 
void foo2() { }
}
#pragma GCC visibility push(default)
__global__ void foo3() { }
#pragma GCC visibility pop

请注意， GCC / Clang 编译器标志 -fvisibility 不会影响这些符号，因为 CUDA 编译器会在发送给主机编译器的代码中使用 attribute((visibility(“hidden”)) 显式标注声明，除非使用了前面描述的一种选择退出机制。

选择加入 ( CUDA 13.0 之前版本)

标记 --device-entity-has-hidden-visibility=true 可以从 CUDA 12.8 开始指定，并且是 CUDA 13.0 及更高版本中的默认值。这将强制隐藏 __global__ 函数和 __managed__/__device__/__constant__ 变量的 ELF 可见性，除非已使用之前指定的选择退出机制之一。

NVCC 变化# 2：强制内部链接

在 CUDA 编程模型中，可以从主机代码启动 __global__ 函数。在发送给主机编译器的代码中，NVCC 会将原始的 __global__ 函数替换为 stub 函数，其中包含对 CUDART 的调用，以在 GPU 上启动内核。为 __managed__/__device__/__constant__ 变量生成类似的存根。

在 NVCC 默认的整个程序编译模式 (-rdc=false) 中，每个翻译单元都会生成一个单独的设备程序。在 CUDA 13.0 之前的工具包中，NVCC 编译器将强制将 __managed__/__device__/__constant__ 存根变量链接到内部链接，但保留了与 __global__ 函数对应的存根的原始链接。

template <typename T>
__global__ void foo() { }  // "foo" stub has external linkage in 
                           // host side code

__managed__ int qqq;       // "qqq" stub has internal linkage in
                           // host side code

带示例的问题概述

许多 CUDA 库 (例如 Thrust) 仅包含报文头，并且报文头中包含 __global__ 模板。如果两个不同的 CUDA 文件 (a.cu 和 b.cu) 包含相同的头文件，并且在整个程序模式 (-rdc=false) 下编译，则每个文件对应的设备程序将完全不同。但是，主机链接器将结合 __global__ 存根函数来自 a.o 和 b.o，因为它们具有外部 (弱) 链接。

这可能会导致意外的运行时行为。当 a.o 和 b.o 在同一程序或库中进行静态链接，或动态链接到单独的共享库并由同一程序加载时，就会出现此问题。

示例 1

//common.h
#include <cassert>
__managed__ int result;
template <typename T>
__global__ void foo() { result = 1; }

//a.cu
#include "common.h"
int first() { 
  foo<int><<<1,1>>>();       // ERROR: may incorrectly launch 
                             // foo<int> in device program created
                             // from b.cu!
  cudaDeviceSynchronize(); 
  return result; 
}

//b.cu
#include "common.h"
int first();
int main() {
  int val = first();
  assert(val == 1);           // assert may fail!
  foo<int><<<1,1>>>();        // ERROR: may incorrectly launch
                              // foo<int> in device program 
                              // created from a.cu!
  cudaDeviceSynchronize();
  assert(result == 1);        // assert may fail!
}

构建为：

$ nvcc  a.cu b.cu -o prog

此时，a.cu 和 b.cu 均需启动内核 foo<int>。根据每个文件创建目标文件 a.o 和 b.o，两个目标文件都包含 foo<int> 的主机侧存根函数。遗憾的是，这两个存根 foo<int> 都具有 (弱) 外部链接，因此主机链接器会合并这些符号，并在最终链接程序中仅选择其中一个符号。

因此，从 a.cu 启动 foo<int> 可能会在 b.o 的设备程序中意外启动 foo<int>，反之亦然。启动的 kernel 可能会更新当前模块中不可见的对象，从而导致意外的 runtime 失败 (例如，a.o 和 b.o 中有不同的 result 副本) 。

$./prog
prog: b.cu:11: int main(): Assertion `result == 1' failed.
Aborted (core dumped)}

请注意，在 CUDA 13.0 及更高版本中，CUDART 将忽略相同主机符号 foo<int> 映射到不同设备程序的重复注册调用。

此问题很难检测和调试，因为没有构建警告或 CUDA 运行时错误，并且有问题的 __global__ 模板可能位于仅第三方头文件库的内部，而这些库的实现对用户来说并不熟悉。

受影响的平台

NVCC 行为的这种变化会影响所有平台，但只能在整个程序编译模式 (-rdc=false) (NVCC 默认模式) 下进行编译。

CUDA 13.0 中引入的解决方案

为避免用户遇到困难，从 CUDA 13.0 开始，NVCC 将强制对发送到主机编译器的代码中生成的 __global__ 函数模板存根函数定义进行内部链接。

在上一个示例中，目标文件 a.o 和 b.o 中的存根函数 foo<int> 将具有内部链接，因此主机链接器在链接期间不会组合这些函数。因此，从 a.cu 启动 foo<int> 现在将在从 a.cu 创建的设备程序中正确启动 foo<int> 内核（对于 b.cu 中的 foo<int> 也是如此）。

选择退出 ( CUDA 13.0)

CUDA 13.0 中引入的更改将破坏一些合法的现有程序。例如，程序中 __global__ 函数模板在一个翻译单元中显式实例化并从另一个翻译单元引用的程序：

//first.cu
template <typename T>
__global__ void foo() { }

template
__global__ void foo<int>(); // explicit instantiation

// second.cu
template <typename T>
__global__ void foo(); // explicit instantiation in first.cu

int main() { foo<int><<<1,1>>>(); cudaDeviceSynchronize(); }

此代码将不再连接 CUDA 13.0 及更高版本，因为在编译 first.o 时，主机存根 foo<int> 的链接将强制进行内部链接，因此主机链接器将无法解析 second.o 中对 foo<int> 的引用。

$nvcc first.cu second.cu -o prog
/usr/bin/ld: /tmp/tmpxft_0032b262_00000000-18_second.o: in function `main':
tmpxft_0032b262_00000000-10_second.cudafe1.cpp:(.text+0xdb): undefined reference to `void foo<int>()'

使用标记 -static-global-template-stub=false 将恢复旧版 NVCC 行为，从而允许构建上述程序。此标志从 CUDA 12.8 开始可用，但在 CUDA 13.0 中将默认值切换为 true。

选择加入 ( CUDA 13.0 之前版本)

要选择加入，请添加 NVCC 标志 -static-global-template-stub=true (自 CUDA 12.8 起提供) 。

总结

CUDA 12.8 及更高版本的 NVCC 标志为您提供了避免一些长期存在的细微运行时错误的工具。考虑在代码中使用它们。请注意，CUDA 13.0 中的标志默认设置会发生变化。这可能会导致某些现有 CUDA C++ 代码的编译或链接失败。考虑更新代码，或使用新标志明确选择退出默认 NVCC 行为。

致谢

感谢以下 NVIDIA 贡献者：Chu-Cheow Lim、Jonathan Bentz 和 Tony Scudiero。