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.cu
extern __global__ void AAA();

int main() { AAA<<<1,1>>>(); }

代碼鏈接為：

$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<<<1,1>>>(); }
	
int main() {

// compiler will mark all instances of foo template as referenced
// from host code, including "foo", which is only actually 
// referenced from device code
foo<<<1,1>>>();
}

• __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) { }
}

template
void 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<<<1,1>>>(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 memory

Incorrect 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 once

Uses 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 177
void foo()
{
  int xxx=0;
}

雜注nv_diagnostic推送和nv_diagnostic彈出可用于保存和恢復當前診斷pragma狀態：

#pragma nv_diagnostic push
#pragma nv_diag_suppress 177
void foo()
{
  int xxx=0;
}
#pragma nv_diagnostic pop
void 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.ptx
13c13
< .func _ZN57_INTERNAL_39_tmpxft_00000a46_00000000_7_test_cpp1_ii_main3fooEv
---
> .func _ZN57_INTERNAL_39_tmpxft_00000a4e_00000000_7_test_cpp1_ii_main3fooEv
16c16
< .visible .global .align 8 .u64 fptr = _ZN57_INTERNAL_39_tmpxft_00000a46_00000000_7_test_cpp1_ii_main3fooEv;
---
> .visible .global .align 8 .u64 fptr = _ZN57_INTERNAL_39_tmpxft_00000a4e_00000000_7_test_cpp1_ii_main3fooEv;
18c18
< .func _ZN57_INTERNAL_39_tmpxft_00000a46_00000000_7_test_cpp1_ii_main3fooEv()
---
> .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 的高級編譯經理。他擁有印度浦那大學計算機工程學士學位。

審核編輯：郭婷