CUDA 11 . 2 的特點是在 GPU 加速應用程序中為設備代碼提供強大的鏈路時間優化（ LTO ）功能。 Device LTO 將設備代碼優化的性能優勢（只有在 nvcc 整個程序編譯模式下才可能）帶到了 CUDA 5 . 0 中引入的 nvcc 單獨編譯模式。

單獨編譯模式允許 CUDA 設備內核代碼跨多個源文件，而在整個程序編譯模式下，程序中的所有 CUDA 設備內核代碼都必須位于單個源文件中。獨立編譯模式將源代碼模塊化引入設備內核代碼，因此是提高開發人員生產率的重要步驟。獨立的編譯模式使開發人員能夠更好地設計和組織設備內核代碼，并使 GPU 加速許多現有的應用程序，而無需進行大量的代碼重構工作，即可將所有設備內核代碼移動到單個源文件中。它還提高了大型并行應用程序開發的開發人員的生產效率，只需要重新編譯帶有增量更改的設備源文件。

CUDA 編譯器優化的范圍通常限于正在編譯的每個源文件。在單獨的編譯模式下，編譯時優化的范圍可能會受到限制，因為編譯器無法看到源文件之外引用的任何設備代碼，因為編譯器無法利用跨越文件邊界的優化機會。

相比之下，在整個程序編譯模式下，程序中存在的所有設備內核代碼都位于同一源文件中，消除了任何外部依賴關系，并允許編譯器執行在單獨編譯模式下不可能執行的優化。因此，在整個程序編譯模式下編譯的程序通常比在單獨編譯模式下編譯的程序性能更好。

使用 CUDA 11 . 0 中預覽的設備鏈接時間優化（ LTO ），可以獲得單獨編譯的源代碼模塊化以及設備代碼整個程序編譯的運行時性能。雖然編譯器在優化單獨編譯的 CUDA 源文件時可能無法進行全局優化的代碼轉換，但鏈接器更適合這樣做。

與編譯器相比，鏈接器具有正在構建的可執行文件的整個程序視圖，包括來自多個源文件和庫的源代碼和符號?？蓤绦形募恼麄€程序視圖使鏈接器能夠選擇最適合單獨編譯的程序的性能優化。此設備鏈接時間優化由鏈接器執行，是 CUDA 11 . 2 中 nvlink 實用程序的一個功能。具有多個源文件和庫的應用程序現在可以通過 GPU 進行加速，而不會影響單獨編譯模式下的性能。

圖 1 。不同編程模式下編譯時和鏈接時優化過程的比較。

圖 1 ，在 nvcc 全程序編譯模式下，要在單個源文件 X . cu 中編譯的設備程序，沒有任何未解析的外部設備函數或變量引用，可以在編譯時由編譯器完全優化。然而，在單獨的編譯模式下，編譯器只能優化正在編譯的單個源文件中的設備代碼，而最終的可執行文件沒有盡可能優化，因為編譯器無法執行跨源文件的更多優化。設備鏈接時間優化通過將優化推遲到鏈接步驟來彌補這一差距。

在設備 LTO 模式下，我們為每個翻譯單元存儲代碼的高級中間形式，然后在鏈接時合并所有這些中間形式以創建所有設備代碼的高級表示。這使鏈接器能夠執行高級優化，例如跨文件邊界內聯，這不僅消除了調用約定的開銷，還進一步支持對內聯代碼塊本身進行其他優化。鏈接器還可以利用已完成的偏移量。例如，共享內存分配是最終確定的，并且數據偏移量僅在鏈路時間已知，因此設備鏈路時間優化現在可以使諸如設備代碼的恒定傳播或折疊之類的低級優化成為可能。即使函數沒有內聯，鏈接器仍然可以看到調用的兩面，以優化調用約定。因此，可以通過設備鏈路時間優化來提高為單獨編譯的程序生成的代碼的質量，并且其性能與以整個程序模式編譯的程序一樣。

為了了解單獨編譯的局限性以及設備 LTO 可能帶來的性能提升，讓我們看一個 MonteCarlo 基準測試中的示例

I 在下面的示例代碼中， MC_Location:: get_domain 不是在另一個文件中定義的標準編譯模式中內聯的，而是
使用 CUDA 11 . 2 中的設備鏈路優化內聯

 __device__ void MCT_Reflect_Particle(MonteCarlo *monteCarlo,
 MC_Particle &particle){

 MC_Location location = particle.Get_Location();
 const MC_Domain &domain = location.get_domain(monteCarlo);
 ...
 ...
 /* uses domain */
 }

函數 get \ u domain 是另一個類的一部分，因此在另一個文件中定義它是有意義的。但是在單獨的編譯模式下，編譯器在調用 get \ u domain （）時將不知道它做什么，甚至不知道它存在于何處，因此編譯器無法內聯該函數，必須隨參數一起發出調用并返回處理，同時也節省空間的事情，如回郵地址后，呼吁。這又使得它無法潛在地優化使用域值的后續語句。在設備 LTO 模式下， get \ u domain （）可以完全內聯，編譯器可以執行更多優化，從而消除調用約定的代碼，并啟用基于域值的優化。

簡而言之，設備 LTO 將所有性能優化都引入到單獨的編譯模式中，而以前只有在整個程序編譯模式中才可用。

使用設備 LTO

要使用設備 LTO ，請將選項 -dlto 添加到編譯和鏈接命令中，如下所示。從這兩個步驟中跳過 -dlto 選項會影響結果。

使用-dlto選項編譯 CUDA 源文件：

nvcc -dc -dlto *.cu

使用-dlto選項鏈接 CUDA 對象文件：

 nvcc -dlto *.o

在編譯時使用 -dlto 選項指示編譯器將正在編譯的設備代碼的高級中間表示（ NVVM-IR ）存儲到 fatbinary 中。在鏈接時使用 -dlto 選項將指示鏈接器從所有鏈接對象檢索 NVVM IR ，并將它們合并到一個 IR 中并執行優化在生成的 IR 上生成代碼。設備 LTO 與任何支持的 SM 架構目標一起工作。

對現有庫使用設備 LTO

設備 LTO 只有在編譯和鏈接步驟都使用 -dlto 時才能生效。如果 -dlto 在編譯時使用，而不是在鏈接時使用，則在鏈接時每個對象都被單獨編譯到 SASS ，然后作為正常鏈接，沒有任何優化機會。如果 -dlto 在鏈接時使用，而不是在編譯時使用，然后鏈接器找不到要執行 LTO 的中間表示，并跳過直接鏈接對象的優化步驟。

如果包含設備代碼的所有對象都是用 -dlto 構建的，那么 Device LTO 工作得最好。但是，即使只有一些對象使用 -dlto ，它仍然可以使用，如圖 2 所示。

圖 2 ：使用非 LTO 庫進行單獨編譯和設備鏈接時間優化。

在這種情況下，在鏈接時，使用 -dlto 構建的對象鏈接在一起形成一個可重定位對象，然后與其他非 LTO 對象鏈接。這不會提供最佳性能，但仍然可以通過在 LTO 對象內進行優化來提高性能。此功能允許使用 -dlto ，即使外部庫不是用 -dlto 構建的；這只是意味著庫代碼不能從設備 LTO 中獲益。

每體系結構的細粒度設備鏈路優化支持

全局 -dlto 選項適用于編譯單個目標體系結構。

使用 -gencode 為多個體系結構編譯時，請確切指定要存儲到 fat 二進制文件中的中間產物。例如，要在可執行文件中存儲 Volta SASS 和 Ampere PTX ，您當前可以使用以下選項進行編譯：

nvcc -gencode arch=compute_70,code=sm_70
 -gencode arch=compute_80,code=compute_80

使用一個新的代碼目標 lto_70 ，您可以獲得細粒度的控制，以指示哪個目標體系結構應該存儲 LTO 中介體，而不是 SASS 或 PTX 。例如，要存儲 Volta LTO 和 Ampere PTX ，可以使用以下代碼示例進行編譯：

nvcc -gencode arch=compute_70,code=lto_70
 -gencode arch=compute_80,code=compute_80

績效結果

設備 LTO 會對性能產生什么樣的影響？

gpu 對內存流量和寄存器壓力非常敏感。因此，設備優化通常比相應的主機優化影響更大。正如預期的那樣，我們觀察到許多應用受益于設備 LTO 。通常，通過設備 LTO 的加速比取決于 CUDA 應用特性。

圖 3 和圖 4 顯示了一個內部基準應用程序和另一個實際應用程序的運行時性能和構建時間的比較圖，這兩個應用程序都采用三種編譯模式：

• 全程序編譯
• 不帶設備 LTO 的單獨編譯
• 使用設備 LTO 模式單獨編譯

我們測試的客戶應用程序有一個占運行時 80% 以上的主計算內核，它調用了分布在不同翻譯單元或源文件中的數百個獨立設備函數。函數的手動內聯是有效的，但如果您希望使用單獨的編譯來維護傳統的開發工作流和庫邊界，則會很麻煩。在這些情況下，使用設備 LTO 來實現潛在的性能優勢而不需要額外的開發工作是非常有吸引力的。

圖 3 ：設備鏈接時間優化的性能加速比優于單獨編譯模式，在某些情況下與整個程序編譯模式相當（越高越好）

如圖 3 所示，帶有設備 LTO 的基準測試和客戶應用程序的運行時性能接近于整個程序編譯模式，克服了單獨編譯模式帶來的限制。請記住，性能的提高在很大程度上取決于應用程序本身的構建方式。正如我們所觀察到的，在某些情況下，收益微乎其微。使用另一個 CUDA 應用程序套件，設備 LTO 的運行時性能平均提高了 25% 左右。

在這篇文章的后面，我們將介紹更多關于設備 LTO 不是特別有用的場景。

除了 GPU 性能之外，設備 LTO 還有另一個方面，那就是構建時間。使用設備 LTO 的總構建時間在很大程度上取決于應用程序大小和其他系統因素。在圖 4 中，內部基準構建時間的相對差異與前面三種不同編譯模式的客戶應用程序進行了比較。內部基準由大約 12000 行代碼組成，而客戶應用程序有上萬行代碼。


有些情況下，由于編譯和優化這些程序所需的過程較少，因此整個程序模式的編譯速度可能更快。此外，在全程序模式下，較小的程序有時可能編譯得更快，因為它有較少的編譯命令，因此對宿主編譯器的調用也較少。但是在全程序模式下的大型程序會帶來更高的優化成本和內存使用。在這種情況下，使用單獨的編譯模式進行編譯會更快。對于圖 4 中的內部基準可以觀察到這一點，其中整個程序模式的編譯時間快了 17% ，而對于客戶應用程序，整個程序模式的編譯速度慢了 25% 。

有限的優化范圍和較小的翻譯單元使單獨編譯模式下的編譯速度更快。當增量更改被隔離到幾個源文件時，單獨的編譯模式還減少了總體的增量構建時間。當啟用設備鏈接時間優化時，編譯器優化階段將被取消，從而顯著減少編譯時間，從而進一步加快單獨編譯模式的編譯速度。但是，同時，由于設備代碼優化階段推遲到鏈接器，并且由于鏈接器可以在單獨編譯模式下執行更多優化，因此單獨編譯的程序的鏈接時間可能隨著設備鏈接時間優化而更高。在圖 4 中，我們可以觀察到設備 LTO 構建時間與基準相比只慢了 7% ，但是與客戶應用程序相比，構建時間慢了近 50% 。

圖 4 ：構建時間加速可以變化（越高越好）。

在 11 . 2 中，我們還引入了新的 nvcc -threads 選項，它在針對多個體系結構時支持并行編譯。這有助于減少構建時間。一般來說，這些編譯模式的總（編譯和鏈接）構建時間可能會因一組不同的因素而有所不同。盡管如此，由于使用設備 LTO 可以顯著縮短編譯時間，我們希望啟用設備鏈接時間優化的單獨編譯模式的總體構建在大多數典型場景中應該是可比的。

設備 LTO 的限制

設備 LTO 在跨文件對象內聯設備功能時特別強大。但是，在某些應用程序中，設備代碼可能都駐留在源文件中，在這種情況下，設備 LTO 沒有太大的區別。

來自函數指針的間接調用（如回調）不會從 LTO 中獲得太多好處，因為這些間接調用不能內聯。

請注意，設備 LTO 執行激進的代碼優化，因此它與使用 -G NVCC 命令行選項來啟用設備代碼的符號調試支持不兼容。

對于 CUDA 11 . 2 ，設備 LTO 只能脫機編譯。設備 LTO 中間窗體尚不支持 JIT LTO 。

像 -maxrregcount 或 -use_fast_math 這樣的文件作用域命令與設備 LTO 不兼容，因為 LTO 優化跨越了文件邊界。如果所有的文件都是用相同的選項編譯的，那么一切都很好，但是如果它們不同，那么設備 LTO 會在鏈接時抱怨。通過在鏈接時指定 -maxrregcount 或 -use_fast_math ，可以覆蓋設備 LTO 的這些編譯屬性，然后該值將用于所有 LTO 對象。

盡管使用設備 LTO 將編譯時優化所花的大部分時間轉移到了鏈接時，但總體構建時間通常在 LTO 構建和非 LTO 構建之間是相當的，因為編譯時間顯著縮短。但是，它增加了鏈接時所需的內存量。我們認為，設備 LTO 的好處應該抵消最常見情況下的限制。

試用設備 LTO

如果您希望在不影響性能或設備源代碼模塊化的情況下，以單獨的編譯模式構建 GPU 加速的應用程序，那么設備 LTO 就適合您了！

使用以單獨編譯模式編譯的設備 LTO 程序可以利用跨文件邊界的代碼優化的性能優勢，從而有助于縮小相對于整個程序編譯模式的性能差距。

為了評估和利用設備 LTO 對 CUDA 應用程序的好處， 立即下載 CUDA 11 . 2 工具包 并進行試用。另外，請告訴我們您的想法。我們一直在尋找改進 CUDA 應用程序開發和運行時性能調優體驗的方法。

關于作者

Mike Murphy 是 NVIDIA 的高級編譯工程師。

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