11 . 2 CUDA C ++編譯器結合了旨在提高開發者生產力和 GPU 加速應用性能的特性和增強。

編譯器工具鏈將 LLVM 升級到 7 . 0 ，這將啟用新功能并有助于改進 NVIDIA GPU 的編譯器代碼生成。設備代碼的鏈接時間優化（ LTO ）（也稱為設備 LTO ）在 CUDA 11 . 0 工具包版本中作為預覽功能引入，現在作為全功能優化功能提供。 11 . 2 CUDA C ++編譯器可以可選地生成一個函數，用于為設備的功能內聯診斷報告，它可以提供編譯器的內聯決策的洞察力。這些診斷報告可以幫助高級 CUDA 開發人員進行應用程序性能分析和調優工作。

CUDA C ++編譯器默認地將設備函數內嵌到調用站點。這使得優化設備代碼的匯編級調試成為一項困難的任務。對于使用 11 . 2 CUDA C ++編譯器工具鏈編譯的源代碼，[EZX223]和 NVIEW 計算調試器可以在調用堆?；厮葜酗@示內聯設備功能的名稱，從而改進調試體驗。

這些和其他新特性被納入 CUDA C ++ 11 . 2 編譯器，我們將在這個帖子中進行深入的跳水。繼續讀！

使用設備 LTO 加速應用程序性能

CUDA 11 . 2 的特點是 設備 LTO ，它為以單獨編譯模式編譯的設備代碼帶來了 LTO 的性能優勢。在 CUDA 5 . 0 中， NVIDIA 引入了獨立編譯模式，以提高開發人員設計和構建 GPU 加速應用程序的效率。沒有單獨的編譯模式，編譯器只支持整個程序編譯模式， CUDA 應用程序中的所有設備代碼必須限制在單個翻譯單元中。單獨的編譯模式使您可以自由地跨多個文件構造設備代碼，包括 GPU 加速的庫和利用增量構建。單獨的編譯模式允許您關注源代碼模塊化。

但是，單獨的編譯模式限制在編譯時可以執行的性能優化范圍內。諸如跨單個翻譯單元邊界的設備函數內聯之類的優化不能在單獨的編譯模式下執行。與整個程序編譯模式相比，這會導致在單獨編譯模式下生成次優代碼，尤其是在針對設備代碼庫進行鏈接時。使用設備 LTO ，在單獨編譯模式下編譯的應用程序的性能與整個編譯模式相當。

LTO 是 CPU 編譯器工具鏈中一個強大的優化功能，我們現在正在使 GPU 加速代碼可以訪問它。對于單獨編譯的設備代碼， Device LTO 支持僅在 NVCC 整個程序編譯模式下才可能進行的設備代碼優化。使用設備 LTO ，您可以利用源代碼模塊化的好處，而不必犧牲整個程序編譯的運行時性能好處。

優化設備代碼的增強調試

我們做了一些增強，以便在某些情況下更容易調試優化的設備代碼。

精確調試

使用 CUDA 11 . 2 ，大多數內聯函數都可以在cuda-gdb和 Nsight 調試器的調用堆棧回溯中看到。您擁有性能優化代碼路徑的一致回溯，更重要的是，您可以更精確地確定錯誤或異常的調用路徑，即使所有函數都是內聯的。

圖 1 顯示了一個場景示例，在調試異常時，此功能可以節省大量時間。

圖 1 。在第 71 行強制數組越界異常的示例代碼

在圖 1 中，函數ExpWrapper調用ForceBoundsException，該函數注入一個數組越界異常。因為函數ForceBoundsException與函數ExpWrapper定義在同一個文件中，所以它只是簡單地內聯在那里。如果沒有對 CUDA 11 . 2 中添加的內聯函數的回溯支持，調用堆棧將只顯示未內聯在此調用路徑中的頂級調用方。在本例中，它恰好是函數ExpWrapper的調用者，因此異常點處的調用堆棧如圖 2 所示，排除了所有其他內聯函數調用。

圖 2 。 CUDA 11 . 2 之前沒有內聯函數的調用堆棧報告行號，沒有完全回溯。

從圖 2 中的調用堆棧可以明顯看出，調用堆棧中的信息非常少，無法有意義地調試最終導致異常點的執行路徑。如果不知道函數是如何內聯的，調用堆棧中提供的行號 71 也沒有用處。在一個三層的深層函數調用中，這個問題看起來很容易找到。隨著堆棧越來越深，這個問題可能會迅速升級。我們知道，這可能是相當令人沮喪的。

圖 3 。在 CUDA 11 . 2 中，一種帶有內聯函數的調用堆棧。

在 CUDA 11 . 2 中， NVIDIA 通過為內聯函數添加有意義的調試信息，朝著優化代碼的符號調試邁出了一步?，F在生成的調用堆棧既精確又有用，包括在每個級別調用的所有函數，包括那些內聯的函數。這使您不僅可以確定發生異常的確切函數，還可以消除觸發異常的確切調用路徑的歧義。

它變得更好了！

更多的調試信息，即使是最優化的代碼

對內聯函數調試的改進不僅是在調用堆棧回溯上查看內聯函數，而且還擴展到源代碼查看。在 CUDA 11 . 2 之前，當函數調用被積極內聯時，反匯編代碼的源代碼視圖是神秘而緊湊的（圖 4 ）。

圖 4 。 CUDA 11 . 2 之前的源代碼反匯編視圖

圖 5 。 CUDA 11 . 2 上啟用源代碼的反匯編代碼視圖。

有更多的調試信息，包括行信息和源代碼行被標記到反匯編代碼段。

圖 5 顯示了 CUDA 11 . 2 上相同反匯編代碼段的源代碼視圖。您可以為優化的代碼段獲得更詳細的源代碼視圖，并且可以單步執行它們。行信息和源代碼行被標記到反匯編源代碼視圖中，即使對于內聯代碼段也是如此。

要啟用此功能，將--generate-line-info（或-lineinfo）選項傳遞給編譯器就足夠了。對優化的設備代碼進行全面的符號調試還不可用。在某些情況下，您可能仍然需要使用-G選項進行調試。然而，僅僅擁有一個精確的調用堆棧和一個詳細的源代碼查看就可以決定性地提高調試性能優化代碼的效率，從而提高開發人員的工作效率。

但還不止這些！

對診斷報告內聯的見解

傳統上，當編譯器做出應用程序開發人員看不到的基于啟發式的優化決策時，編譯器有點像黑匣子。

其中一個關鍵的優化就是函數內聯。如果沒有對匯編輸出進行繁重的后處理，就很難理解內聯的編譯器啟發式方法。只要知道哪些函數是內聯的，哪些不是內聯的，就可以節省很多時間，這就是我們在 CUDA 11 . 2 中介紹的?，F在您不僅知道函數何時沒有內聯，而且還知道為什么函數不能內聯。然后可以重構代碼，向函數 de Clara 選項添加內聯關鍵字，或者執行其他源代碼重構（如果可能的話）。

您可以通過一個新選項--optimization-info=inline獲得關于優化器內聯決策的診斷報告。啟用內聯診斷時，當函數無法內聯時，優化器會報告其他診斷。

圖 6 。樣品測試. cu 用于以下內聯診斷生成的文件。

remark: test.cu:16:12: _Z7callee2i inlined into _Z6callerPii with cost=always

remark: test.cu:17:11: _Z7callee3i inlined into _Z6callerPii with cost=always

remark: test.cu:18:12: _Z7callee1i not inlined into _Z6callerPii because it should never be inlined (cost=never)

在某些情況下，您可能會得到更詳細的診斷：

remark: x.cu:312:28: callee not inlined into caller because callee doesn't have forceinline attribute and is too big for auto inlining (CalleeSize=666)

有關內聯的診斷報告對于重構代碼以適當地使用內聯函數的性能優勢非常有用。內聯診斷在編譯器運行內聯過程時發出。當從編譯器多次調用內聯程序時，前一個過程中未內聯的調用站點可能會內聯到后一個過程中通過。那個 CUDA C ++編譯器文檔解釋了如何在 NVCC 調用期間使用此選項。

通過并行編譯減少構建時間

可以使用 -gencode/-arch/-code 命令行選項同時調用 CUDA C ++編譯器，以編譯多個 GPU 架構的 CUDA 設備代碼。雖然這是一個方便的特性，但它可能會導致由于幾個中間步驟而增加構建時間。

特別地，編譯器需要對 CUDA C ++源代碼進行多次處理，并使用不同的 __CUDA__ARCH__ 內置宏的值來指定每個不同的計算架構，包括額外的預處理步驟，其中內置的宏未被定義，以編譯主機平臺的源代碼。之后，預處理的 CUDA C ++設備代碼實例必須編譯成指定的每個目標 GPU 架構的機器代碼。這些步驟目前是連續進行的。

為了減輕由多個編譯過程產生的編譯時間的增加，從 CUDA 11 。 2 版本開始， CUDA C ++編譯器支持一個新的 —threads 《number》 命令行選項（簡稱-t）來生成單獨的線程以并行執行獨立編譯傳遞。如果在單個 nvcc 命令中編譯多個文件， -t 將并行編譯這些文件。 參數確定 NVCC 編譯器為并行執行獨立編譯步驟而生成的獨立輔助線程數。

對于特殊情況 -t0 ，使用的線程數是機器上的 CPU 數。當調用 NVCC 為多個 GPU 架構同時編譯 CUDA 設備代碼時，此選項有助于減少總體構建時間。默認情況下，這些步驟是連續執行的。

Example

以下命令為兩個虛擬體系結構生成。 ptx 文件： compute_52 和 compute_70 。對于 compute_52 ，為兩個 GPU 目標生成。 cubin 文件： sm_52 和 sm_60 ；對于 compute_70 ，為 sm_70. 生成。 cubin 文件

nvcc -gencode arch=compute_52,code=sm_52 -gencode arch=compute_52,code=sm_60 -gencode arch=compute_70,code=sm_70 t.cu

并行編譯有助于在編譯大量應用 CUDA C ++設備代碼到多個 GPU 目標的應用程序時減少總體構建時間。如果源代碼主要是 C / C ++主機代碼，只有少量 CUDA 設備代碼，或者如果僅以單個虛擬架構/ GPU-SM 組合為目標，則可能不會減少整個構建時間。換句話說，構建時的加速可能會因程序、編譯目標特性以及 NVCC 可以生成的并行編譯線程的數量而異。

NVCC 啟動 helper 線程來動態地并行執行編譯步驟（如 CUDA 編譯軌跡圖 中所描述的），受編譯步驟之間的序列化依賴關系的約束，其中編譯步驟僅在其依賴的所有先前步驟完成之后才在單獨的線程上啟動。

圖 7 顯示了當 NVCC 生成的獨立編譯線程的限制增加時（ -t N 選項），由于并行編譯而導致的 CUDA 編譯加速是如何變化的。這適用于需要不同級別的獨立編譯步驟的編譯軌跡，這些步驟可以并行執行。

圖 7 。為多個 GPU 架構編譯 NVIDIA 性能原語（ NPP ）的并行編譯加速。

CPU 型號： i7-7800X CPU @ 3 。 50GHz # CPU ： 12 ，每核線程數： 2 ，每插槽核數： 6 ，內存： 31G 。 （所有的編譯都使用 make-j8 ）

NVCC 并行線程編譯特性可以與進程級構建并行性（即， make -j N ）一起使用。但是，必須考慮主機平臺的特性，以避免過度訂閱生成系統資源（例如， CPU 核心數、可用內存、其他工作負載），這可能會對總體生成時間產生負面影響。

新的編譯器內置提示，可以更好地優化設備代碼

CUDA 11 。 2 支持新的內置程序，使您能夠向編譯器指示編程提示，以便更好地生成和優化設備代碼。

使用 __builtin_assume_aligned ， 可以向編譯器提示指針對齊，編譯器可以使用指針對齊進行優化。類似地， __builtin_assume 和 __assume 內置可以用來指示運行時條件，以幫助編譯器生成更好的優化代碼。下一節將深入研究每個特定的內置提示函數。

void * __builtin_assume_aligned(const void *ptr, size_t align)
void *__builtin_assume_aligned(const void *ptr, size_t align, offset)

__builtin_assume_aligned內置函數可用于向編譯器提示作為指針傳遞的參數至少與align字節對齊。當參數(char *)ptr - offset至少與align字節對齊時，可以使用帶有offset的版本。兩個函數都返回參數指針。

編譯器可以使用這種對齊提示來執行某些代碼優化，如加載/存儲矢量化，以更好地工作?？紤]一下這里顯示的函數中的示例代碼，該函數使用內置函數來指示參數ptr可以假定至少與 16 個字節對齊。

__device int __get(int*ptr)
{
int *v = static_cast 
(__builtin_assume_aligned(ptr, 16));
return *v + *(v+1) + *(v+2) + *(v+3);
}

前面的代碼示例在使用nvcc -rdc=true -ptx foo.cu編譯時沒有內置函數，生成了以下 PTX ，其中對返回表達式執行了四個單獨的加載操作。

ld.u32 %r1, [%rd1];
ld.u32 %r2, [%rd1 + 4];
ld.u32 %r4, [%rd1 + 8];
ld.u32 %r6, [%rd1 +12];

當使用內置函數向編譯器提示指針是 16 字節對齊的時，生成的 PTX 反映了這樣一個事實：編譯器可以將加載操作組合成一個向量化的加載操作。

ld.v4.u32 {%r1, %r2, %r3, %r4 }, [%rd1];

由于四個加載是并行執行的，因此單個矢量化加載操作所需的執行時間更少。這避免了向內存子系統發出多個請求的開銷，同時還保持了較小的二進制大小。

void * __builtin_assume(bool exp)

__builtin__assume內置函數允許編譯器假定提供的布爾參數為 true 。如果參數在運行時不為 true ，則行為未定義。參數表達式不能有副作用。盡管 CUDA 11 . 2 文檔指出副作用已被丟棄，但此行為在將來的版本中可能會發生更改，因此可移植代碼在提供的表達式中不應產生副作用。

例如，對于下面的代碼段， CUDA 11 . 2toolkit 編譯器可以用更少的指令優化 modulo-16 操作，因為知道num變量的值是肯定的。

__device__ int mod16(int num) 
{
      __builtin_assume(num > 0); 
      return num % 16; 
}

如下一個生成的 PTX 代碼示例所示，當使用nvcc -rdc=true -ptx編譯示例代碼時，編譯器為模運算生成一條 AND 指令。

ld.param.u32   %r1, [_Z5Mod16i_param_0]; 
     and.b32   %r2, %r1, 15;
st.param.b32   [func_retval0+0], %r2;

如果沒有提示，編譯器必須考慮num值為負值的可能性，如生成的 PTX 代碼（包括附加指令）所示。

ld.param.u32   %r1, [_Z5Mod16i_param_0];
     shr.s32   %r2, %r1, 31;
     shr.u32   %r3, %r2, 28;
     add.s32   %r4, %r1, %r3;
     and.b32   %r2, %r1, 15;
     sub.s32   %r6, %r1, %r5
st.param.b32   [func_retval0+0], %r2;

使用時， NVCC 還支持類似的內置函數__assume(bool)cl . exe 文件作為主機編譯器。

void * __builtin_unreachable(void)

在 CUDA 11 . 3 中，我們將介紹__builtin_unreachable內置函數。這個內置函數在 CUDA 11 . 3 中引入時，可用于向編譯器指示控制流永遠不會到達調用此函數的點。如果控制流在運行時到達該點，則程序具有未定義的行為。此提示可以幫助代碼優化器生成更好的代碼：

__device__ int get(int input)
{
   switch (input)
   {
          case 1: return 4;
          case 2: return 10;
          default: __builtin_unreachable();
   }
}

用 CUDA 11 . 3 中的nvcc -rdc=true -ptx編譯早期代碼片段生成的 PTX 將把整個 switch 語句優化為一條 SELECT 指令。

  ld.param.u32   %r1, [_Z3geti_param_0];
   setp.eq.s32   %p1, %r1, 1;
      selp.b32   %r2, 4, 10, %p1;
  st.param.b32   [func_retval0+0], %r2;

如果沒有__builtin_unreachable調用，編譯器將生成一個警告，指出控制流已到達非 void 函數的結尾。通常，必須注入一個偽返回 0 以避免出現警告消息。

__device__ int get(int input)
{
switch (input)
{
      case 1: return 4;
     case 2: return 10;
      default: return 0;
}
}

添加 return 以避免編譯器警告會導致更多的 PTX 指令，這也有抑制進一步優化的潛在副作用。

  ld.param.u32   %r1, [_Z3geti_param_0];
setp.eq.s32%p1, %r1, 2;
selp.b32%r2, 10, 0, %p1;
setp.eq.s32%p2, %r1, 1;
selp.b32%r3, 4, %r2, %p2;
  st.param.b32[func_retval0+0], %r2;

__builtin_assume和__builtin_assume_aligned函數在內部映射到llvm.assumeLLVM 內在函數。

“ 請注意，優化器 MIG ht 限制對 llvm.assume 保留僅用于形成內在函數輸入參數的指令。如果用戶提供的額外信息 llvm.assume 內在的并不能導致代碼質量的全面提高。因此， llvm.assume 不應用于記錄優化器可以以其他方式推斷的基本數學不變量或對優化器沒有多大用處的事實。”

某些主機編譯器可能不支持早期的內置函數。在這種情況下，必須注意在代碼中調用內置函數的位置。

下表給出了主機編譯器為 gcc 時使用 __builtin_assume 的示例。由于 gcc 不支持此內置函數，因此在未定義 __CUDA_ARCH__ 宏的主機編譯階段，對 __builtin_assume 的調用不應出現在 __device__ 函數之外。

表 1 。當主機編譯器不支持內置項時，使用內置項的示例。

有關如何使用這些內置函數的詳細信息，請參閱 編譯器優化提示函數 。

警告可以被抑制或標記為錯誤

NVCC 現在支持可以用來管理編譯器診斷的命令行選項。您可以選擇讓編譯器隨診斷消息一起發出錯誤號，并指定編譯器應將與錯誤號關聯的診斷視為錯誤還是完全抑制。這些選項不適用于主機編譯器或預處理器發出的診斷。在將來的版本中，編譯器還將支持 pragmas ，以將特定的警告提升到錯誤或抑制它們。

Usage

--display-error-number （-err-no）

顯示 CUDA 前端編譯器生成的任何消息的診斷號。

--diag-error 《error-number》，。.. （-diag-error）

為 CUDA 前端編譯器生成的指定診斷消息發出錯誤。

--diag-suppress 《error-number》，。.. （-diag-suppress）

抑制 CUDA 前端編譯器生成的指定診斷消息。

Example

設備函數 hdBar 調用主機函數 hostFoo 并且變量 i 在 hostFoo 中未使用的示例代碼：

void hostFoo(void)
{
int i = 0;
}
__host__ __device__ void hdBar(bool cond)
{
if (cond)
hostFoo();
}

以下代碼示例顯示帶有默認警告的診斷號：

$nvcc -err-no -ptx warn.cu
warn.cu(1): warning #177-D: variable "i" was declared but never referenced
warn.cu(2): warning #20011-D: calling a __host__ function("hostFoo()") from a __host__ __device__ function("hdBar") is not allowed

以下代碼示例將警告# 20011 升級為錯誤：

$nvcc -err-no -ptx -diag-error 20011 warn.cu
warn.cu(1): warning #177-D: variable "i" was declared but never referenced
warn.cu(2): error: calling a __host__ function("hostFoo()") from a __host__ __device__ function("hdBar") is not allowed

以下代碼示例禁止顯示警告# 20011 ：

$nvcc -err-no -ptx -diag-suppress 20011 warn.cu
warn.cu(1): warning #177-D: variable "i" was declared but never referenced

NVVM 升級到 LLVM 7 。 0

CUDA 11 。 2 編譯器工具鏈接收 LLVM7 。 0 升級。

升級到 LLVM 7 。 0 將打開通向此 LLVM 版本中存在的新功能的大門。它通過利用 LLVM 7 中可用的新優化，為進一步實現性能調整工作提供了更堅實的基礎。

圖 8 顯示了使用包含基于 LLVM7 。 0 的高級 NVVM 優化器的 11 。 2 編譯器工具鏈編譯的 HPC 應用程序子集對基于 Volta 和 Ampere 的 GPU 的運行時性能影響，而 11 。 1 編譯器工具鏈包含基于 LLVM3 。 4 的高級 NVVM 優化器。

圖 8 。 HPC 應用程序套件的 Geomean 性能增益/損失

相對于 LLVM3 。 4 ，基于 A100 和 V100 的 NVVM 。

libnvm 升級到 LLVM 7 。 0

使用 CUDA 11 。 2 版本， CUDA C ++編譯器， LIbvvm 和 NVRTC 共享庫都已升級到 LLVM 7 代碼庫。 libNVVM 庫為 LLVM 提供了 GPU 擴展，以支持更廣泛的社區，包括編譯器、 DSL 轉換器和針對 NVIDIA GPU 上計算工作負載的并行應用程序。 NVRTC 共享庫有助于在運行時編譯動態生成的 CUDA C ++源代碼。

由于 libNVVM 庫包含 llvm7 。 0 支持， libnvvmapi 和 nvvmir 規范已修改為與 llvm7 。 0 兼容。要更新輸入 IR 格式，請參閱已發布的 NVVM IR 規范。此 libNVVM 升級與以前版本中支持的調試元數據 IR 不兼容。依賴于調試元數據生成的第三方編譯器應該適應新的規范。在這次升級中， libnvm 也不推薦使用文本 IR 接口。我們建議您使用 LLVM 7 。 0 位碼格式 。有關對基于 libnvm 的編譯器軟件所做更改的更多信息，請參閱 libNVVM 規范 和 NVVM IR 規范 。

此升級還帶來了對源代碼級調試支持的增強。編譯器前端可能需要一個矮型表達式來指示在運行時保存變量值的位置。如果沒有對 DWARF 表達式的適當支持，則無法在調試器中檢查此類變量。 libNVVM 升級的一個重要方面是，使用 DWARF 表達式之類的操作可以更廣泛地表達這些變量位置。 NVVM IR 現在使用 本質與操作 支持此類表達式。這樣一個變量的最終位置用這些表達式用 DWARF 表示

試試 CUDA 11 。 2 編譯器的功能

CUDA 11 。 2 工具包包含了一些專注于提高 GPU 性能和提升開發人員體驗的功能。［VZX107 型］

編譯器工具鏈升級到 LLVM 7 、設備 LTO 支持和新編譯器內置的能力，這些能力可以利用來增強 CUDA C ++應用程序的性能。

對內聯設備函數的虛擬堆?；厮葜С?、關于函數內聯決策的編譯器報告、并行 CUDA 編譯支持以及控制編譯器警告診斷的能力是 CUDA 11 。 2 工具包中的新功能，旨在提高您的生產效率。

關于作者

Arthy Sundaram 是 CUDA 平臺的技術產品經理。她擁有哥倫比亞大學計算機科學碩士學位。她感興趣的領域是操作系統、編譯器和計算機體系結構。

Jaydeep Marathe 是 NVIDIA 的高級編譯工程師。他擁有北卡羅來納州立大學計算機科學碩士和博士學位。

Hari Sandanagobalane 是 NVIDIA 的高級編譯工程師。他擁有新加坡國立大學計算機科學碩士學位。

Gautam Chakrabarti 是 NVIDIA 的高級編譯工程師。他的興趣包括編譯器技術，使高效的程序執行和改善開發人員的經驗。他擁有密歇根州立大學計算機科學碩士學位。

Mukesh Kapoor 是 NVIDIA 的高級編譯工程師。他擁有印度班加羅爾工業學院自動化碩士學位。Steve Wells 是 NVIDIA 的高級調試器工程師。他擁有滑鐵盧大學的數學學士學位。Mike Murphy 是 NVIDIA 的高級編譯工程師。

審核編輯：郭婷