在下一個 CUDA 主要版本 CUDA 13.0 中，NVIDIA 將對 NVIDIA CUDA 編譯器驅動程序 (NVCC) 作出兩項重大更改，這些更改將影響 __global__ 函數和設備變量的 ELF 可見性和關聯。這些更新旨在防止長期以來難以檢測和調試的細微運行時錯誤。但是，這些更改可能會影響一些現有的 CUDA C++ 程序。
本文旨在提醒用戶注意潛在的中斷，解釋更改背后的原因，并就可以恢復舊版行為的 NVCC 標志提供指導。表 1 匯總了這兩項更改。
| 特征 | ELF 可見性 | 強制內部關聯 |
| 特征 詳細信息 | 強制隱藏的 ELF 可見性 for __global__ 函數、__managed__/__device__/__constant__ 變量 |
強制 __global__ 函數主機模板存根定義具有內部鏈接 (僅限整個程序模式) |
| 受影響的平臺 | 非 Windows 上的 Shared libraries | 所有平臺均處于 NVCC 整個程序編譯模式 (-rdc=false) 。這是默認的 NVCC 模式。 |
| 用戶影響 | 默認情況下，__global__ 函數、__managed__/__device__/__shared__ 變量不會從共享庫中導出 |
在另一個翻譯單元中對 __global__ 模板實例化的引用將無法構建。 |
| 控制標志 (CUDA 12.8+) | “-device-entity-has-hidden-visibility={true|false} CUDA 13.0+ 中的默認設置：true CUDA < 13.0 中的默認設置：false” |
-static-global-template-stub={true|false} CUDA 13.0+ 中的默認設置：true CUDA < 13.0 中的默認設置：false |
| 選擇退出 ( CUDA 13.0 及以上) | -device-entity-has-hidden-visibility=false |
-static-global-template-stub=false |
| 選擇加入 (CUDA 12.8+) | -device-entity-has-hidden-visibility=true |
-static-global-template-stub=true |
__global__ 函數和設備變量的 ELF 可見性和關聯 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:44 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 副本) 。
$./progprog: 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.cutemplate <typename T>__global__ void foo() { }template__global__ void foo<int>(); // explicit instantiation |
// second.cutemplate <typename T>__global__ void foo(); // explicit instantiation in first.cuint 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。
?
