使用 clang 編譯 CUDA¶

簡介 ¶

本文檔描述如何使用 clang 編譯 CUDA 程式碼，並提供有關 LLVM 和 clang 的 CUDA 實作的一些細節。

本文檔假設您對 CUDA 有基本的了解。有關 CUDA 程式設計的資訊，請參閱 CUDA 程式設計指南。

編譯 CUDA 程式碼 ¶

先決條件 ¶

自 llvm 3.9 版以來，已支援 CUDA。Clang 目前支援 CUDA 7.0 到 12.1。如果 clang 偵測到更新的 CUDA 版本，它會發出警告，並嘗試使用偵測到的 CUDA SDK，就好像它是 CUDA 12.1 一樣。

在建置 CUDA 程式碼之前，您需要先安裝 CUDA SDK。如需詳細資訊，請參閱 NVIDIA 的 CUDA 安裝指南。請注意，clang 可能不支援某些 Linux 套件管理器安裝的 CUDA 工具包。Clang 確實會嘗試處理一些常見 Linux 發行版上 CUDA 安裝的特定細節，但一般來說，最可靠的方法是從 NVIDIA 的 .run 套件在單一目錄中安裝 CUDA，並透過 –cuda-path=… 參數指定其位置。

Linux 上支援 CUDA 編譯。MacOS 和 Windows 上的編譯可能會也可能不會成功，而且目前沒有維護者。

呼叫 clang ¶

呼叫 clang 進行 CUDA 編譯的方式與編譯一般 C++ 類似。您只需要注意一些額外的旗標。

您可以使用這個程式作為一個簡單的範例。將其儲存為 axpy.cu。（Clang 會透過辨識檔名是否以 .cu 結尾來判斷您是否正在編譯 CUDA 程式碼。或者，您也可以傳遞 -x cuda。）

要建置並執行，請執行以下命令，並按照下方說明填寫括號中的部分

$ clang++ axpy.cu -o axpy --cuda-gpu-arch=<GPU arch> \
    -L<CUDA install path>/<lib64 or lib>             \
    -lcudart_static -ldl -lrt -pthread
$ ./axpy
y[0] = 2
y[1] = 4
y[2] = 6
y[3] = 8

在 MacOS 上，將 -lcudart_static 替換為 -lcudart；否則，當您執行程式時，可能會收到「CUDA 驅動程式版本不足以執行 CUDA 執行階段版本」的錯誤訊息。

<CUDA 安裝路徑> – 您安裝 CUDA SDK 的目錄。通常為 /usr/local/cuda。

如果以 64 位元模式編譯，請傳遞例如 -L/usr/local/cuda/lib64；否則，請傳遞例如 -L/usr/local/cuda/lib。（在 CUDA 中，設備程式碼和主機程式碼始終具有相同的指標寬度，因此如果您要為主機編譯 64 位元程式碼，則也會為設備編譯 64 位元程式碼。）請注意，從 v10.0 CUDA SDK 開始不再支援 32 位元應用程式的編譯。
<GPU 架構> – 您 GPU 的運算能力。例如，如果您要在運算能力為 3.5 的 GPU 上執行程式，請指定 --cuda-gpu-arch=sm_35。

注意：您不能將 compute_XX 作為參數傳遞給 --cuda-gpu-arch；目前僅支援 sm_XX。但是，clang 總是會在其二進制檔案中包含 PTX，因此例如使用 --cuda-gpu-arch=sm_30 編譯的二進制檔案將與例如 sm_35 GPU 向前相容。

您可以多次傳遞 --cuda-gpu-arch 來為多個架構編譯。

-L 和 -l 旗標只需要在連結時傳遞。在編譯時，如果您沒有將 CUDA SDK 安裝到 /usr/local/cuda 或 /usr/local/cuda-X.Y 中，則可能還需要傳遞 --cuda-path=/path/to/cuda。

控制數值程式碼的旗標 ¶

如果您使用的是 GPU，您可能會關心如何讓數值程式碼快速執行。與大多數 CPU 相比，GPU 硬體允許對數值運算進行更多控制，但這也導致您需要處理更多編譯器選項。

您可能希望調整的旗標包括

-ffp-contract={on,off,fast}（在編譯 CUDA 時，主機和設備上的預設值為 fast）控制編譯器是否發出融合乘加運算。
- off：永遠不發出 fma 運算，並防止 ptxas 融合乘法和加法指令。
- on：融合單一語句中的乘法和加法，但絕不跨語句融合（C11 語義）。防止 ptxas 融合其他乘法和加法。
- fast：只要有利可圖，即使跨語句，也會融合乘法和加法。不會阻止 ptxas 融合其他乘法和加法。
融合乘加指令可能比未融合的等效指令快得多，但由於 fma 中的中間結果沒有經過四捨五入，因此此標誌可能會影響數值程式碼。
-fcuda-flush-denormals-to-zero（預設值：關閉）啟用此選項後，浮點運算可能會將非正規數輸入和/或輸出刷新為 0。對非正規數的運算通常比對正規數的相同運算慢得多。
-fcuda-approx-transcendentals（預設值：關閉）啟用此選項後，編譯器可能會發出對更快、近似的超越函數版本的呼叫，而不是使用較慢、完全符合 IEEE 標準的版本。例如，此標誌允許 clang 發出 ptx sin.approx.f32 指令。

這是由 -ffast-math 隱含的。

標準函式庫支援 ¶

在 clang 和 nvcc 中，大部分 C++ 標準函式庫在設備端都不受支援。

`<math.h>` 和 `<cmath>`¶

在 clang 中，math.h 和 cmath 可用，並且通過了改編自 libc++ 測試套件的測試。

在 nvcc 中，math.h 和 cmath 大多可用。命名空間 std 中 ::foof 的版本（例如 std::sinf）不可用，並且如果標準要求採用整數參數的重載，則這些版本通常不可用。

#include <math.h>
#include <cmath.h>

// clang is OK with everything in this function.
__device__ void test() {
  std::sin(0.); // nvcc - ok
  std::sin(0);  // nvcc - error, because no std::sin(int) override is available.
  sin(0);       // nvcc - same as above.

  sinf(0.);       // nvcc - ok
  std::sinf(0.);  // nvcc - no such function
}

`<std::complex>`¶

nvcc 沒有正式支援 std::complex。在 __device__ 程式碼中使用 std::complex 是一個錯誤，但由於 nvcc 對「錯誤端規則」（見下文）的解釋，它在 __host__ __device__ 程式碼中通常可以正常運作。但是，我們從實作者那裡聽說，可能會遇到 nvcc 省略對 std::complex 函數的呼叫的情況，尤其是在沒有優化的情況下編譯時。

從 2016-11-16 開始，clang 支援 std::complex，沒有這些注意事項。它已經過 libstdc++ 4.8.5 及更新版本的測試，但已知僅適用於 2016-11-16 之後的 libc++。

`<algorithm>`¶

在 C++14 中，<algorithm> 中的許多有用函數（特別是 std::min 和 std::max）變成了 constexpr。因此，在使用 clang 編譯時，您可以在設備程式碼中使用這些函數。

從程式碼中偵測 clang 與 NVCC ¶

雖然 clang 的 CUDA 實作與 NVCC 大致相容，但您可能仍然想在使用 clang 編譯 CUDA 程式碼時進行偵測。

這很棘手，因為 NVCC 可能會在其編譯過程中呼叫 clang！例如，NVCC 在編譯設備程式碼時會使用主機編譯器的預處理器，而該主機編譯器實際上可能是 clang。

當 clang 實際編譯 CUDA 程式碼（而不是作為 NVCC 的子工具使用）時，它會定義 __CUDA__ 巨集。 __CUDA_ARCH__ 僅在設備模式下定義（但如果 NVCC 使用 clang 作為預處理器，則會定義）。因此，您可以使用以下咒語來偵測主機和設備模式下的 clang CUDA 編譯：

#if defined(__clang__) && defined(__CUDA__) && !defined(__CUDA_ARCH__)
// clang compiling CUDA code, host mode.
#endif

#if defined(__clang__) && defined(__CUDA__) && defined(__CUDA_ARCH__)
// clang compiling CUDA code, device mode.
#endif

clang 和 nvcc 都在 CUDA 編譯期間定義 __CUDACC__。您可以透過尋找 __NVCC__ 來專門偵測 NVCC。

clang 和 nvcc 之間的方言差異 ¶

沒有正式的 CUDA 規範，clang 和 nvcc 使用的語言方言略有不同。以下，我們描述了一些差異。

本節內容比較枯燥；希望您可以跳過本節，並過上幸福無憂的生活。

編譯模型 ¶

clang 和 nvcc 之間的大多數差異源於 clang 和 nvcc 使用的不同編譯模型。nvcc 使用*分離式編譯*，其工作原理大致如下：

對輸入的 .cu 檔案執行預處理器，將其拆分為兩個原始程式檔：H，包含主機的原始程式碼，以及 D，包含設備的原始程式碼。

對於我們要編譯的每個 GPU 架構 arch，執行以下操作：

使用 nvcc 本身編譯 D。結果是一個針對 P_arch 的 ptx 檔案。

（可選）呼叫 PTX 組譯器 ptxas，為 arch 生成一個包含 GPU 機器碼（SASS）的檔案 S_arch。

呼叫 fatbin 將所有 P_arch 和 S_arch 檔案組合成一個「胖二進位」檔案 F。

使用外部主機編譯器（gcc、clang 或您喜歡的任何編譯器）編譯 H。 F 被打包成一個頭文件，並強制包含到 H 中；nvcc 生成的程式碼會呼叫此頭文件以執行啟動核心等操作。

clang 使用*合併式解析*。這類似於分離式編譯，不同之處在於所有主機和設備程式碼都存在，並且在兩個編譯步驟中都必須在語義上正確。

對於我們要編譯的每個 GPU 架構 arch，執行以下操作：

使用 clang 編譯輸入的 .cu 檔案以供設備使用。即使我們目前沒有為主機產生程式碼，__host__ 程式碼也會被解析，並且必須在語義上正確。

此步驟的輸出是一個 ptx 檔案 P_arch。

呼叫 ptxas 來產生一個 SASS 檔案 S_arch。請注意，與 nvcc 不同，clang 總是產生 SASS 程式碼。

呼叫 fatbin 將所有 P_arch 和 S_arch 檔案組合成一個單一的胖二進制檔案 F。

使用 clang 編譯 H。即使我們目前沒有為設備產生程式碼，__device__ 程式碼也會被解析，並且必須在語義上正確。

F 會被傳遞到這個編譯過程中，clang 會將其包含在一個特殊的 ELF 區段中，以便像 cuobjdump 這樣的工具可以找到它。

（您可能會在這裡問，為什麼 clang 需要多次解析輸入檔案？為什麼不只解析一次，然後使用 AST 為主機和每個設備架構產生程式碼呢？

不幸的是，這是行不通的，因為我們必須在主機編譯和針對每個 GPU 架構的設備編譯期間定義不同的巨集。）

clang 的方法使其對 C++ 的邊緣情況具有高度的穩健性，因為它不需要在早期階段就決定保留哪些聲明和捨棄哪些聲明。但是，它也有一些您應該注意的後果。

基於 `host` 和 `device` 屬性的重載 ¶

讓我們用「H」、「D」和「HD」分別代表「__host__ 函數」、「__device__ 函數」和「__host__ __device__ 函數」。沒有屬性的函數的行為與 H 相同。

nvcc 不允許您創建具有相同簽章的 H 和 D 函數

// nvcc: error - function "foo" has already been defined
__host__ void foo() {}
__device__ void foo() {}

但是，nvcc 允許您「重載」具有不同簽章的 H 和 D 函數

// nvcc: no error
__host__ void foo(int) {}
__device__ void foo() {}

在 clang 中，__host__ 和 __device__ 屬性是函數簽章的一部分，因此允許存在具有（其他方面）相同簽章的 H 和 D 函數

// clang: no error
__host__ void foo() {}
__device__ void foo() {}

HD 函數不能被具有相同簽章的 H 或 D 函數重載

// nvcc: error - function "foo" has already been defined
// clang: error - redefinition of 'foo'
__host__ __device__ void foo() {}
__device__ void foo() {}

// nvcc: no error
// clang: no error
__host__ __device__ void bar(int) {}
__device__ void bar() {}

在解析重載函數時，clang 會考慮呼叫者和被呼叫者的主機/設備屬性。這些屬性在重載解析期間用作決勝因素。有關完整規則集，請參閱 IdentifyCUDAPreference，但概括來說，它們是

D 函數更喜歡呼叫其他 D 函數。HD 函數的優先級較低。

同樣地，H 函數更喜歡呼叫其他 H 函數或 __global__ 函數（優先級相同）。HD 函數的優先級較低。

HD 函數更喜歡呼叫其他 HD 函數。

在編譯裝置程式碼時，HD 函式呼叫 Ds 函式的優先順序會低於 HD 函式，而呼叫 Hs 函式的優先順序又更低。如果必須呼叫 H 函式，則在為此 HD 函式產生程式碼時，程式碼就會發生格式錯誤。我們將此稱為「錯誤端規則」，請參閱以下範例。

在編譯主機程式碼時，這些規則是對稱的。

一些範例

__host__ void foo();
__device__ void foo();

__host__ void bar();
__host__ __device__ void bar();

__host__ void test_host() {
  foo();  // calls H overload
  bar();  // calls H overload
}

__device__ void test_device() {
  foo();  // calls D overload
  bar();  // calls HD overload
}

__host__ __device__ void test_hd() {
  foo();  // calls H overload when compiling for host, otherwise D overload
  bar();  // always calls HD overload
}

錯誤端規則範例

__host__ void host_only();

// We don't codegen inline functions unless they're referenced by a
// non-inline function.  inline_hd1() is called only from the host side, so
// does not generate an error.  inline_hd2() is called from the device side,
// so it generates an error.
inline __host__ __device__ void inline_hd1() { host_only(); }  // no error
inline __host__ __device__ void inline_hd2() { host_only(); }  // error

__host__ void host_fn() { inline_hd1(); }
__device__ void device_fn() { inline_hd2(); }

// This function is not inline, so it's always codegen'ed on both the host
// and the device.  Therefore, it generates an error.
__host__ __device__ void not_inline_hd() { host_only(); }

就錯誤端規則而言，樣板函式的行為也像 inline 函式：除非這些函式被實體化（通常是在呼叫它們的過程中），否則不會產生程式碼。

clang 在錯誤端規則方面的行為與 nvcc 相符，但 nvcc 只會針對 not_inline_hd 發出警告；允許裝置程式碼呼叫 not_inline_hd。在其產生的程式碼中，nvcc 可能會完全省略 not_inline_hd 對 host_only 的呼叫，或者它可能會嘗試為裝置上的 host_only 產生程式碼。您得到的結果似乎取決於編譯器是否選擇內嵌 host_only。

成員函式（包括建構函式）可以使用 H 和 D 屬性進行多載。但是，解構函式不能多載。

Clang 對主機和裝置函式宣告的警告 ¶

當 Clang 偵測到主機 (H) 和裝置 (D) 函式使用相同的簽章宣告或定義時，它會發出警告。預設情況下，這些警告不會啟用。

若要啟用這些警告，請使用以下編譯器旗標

-Wnvcc-compat

在主機/裝置上使用不同的類別 ¶

有時您可能希望擁有一個具有不同主機/裝置版本的類別。

如果類別的所有成員在主機和裝置上都相同，則您只需為類別的成員函式提供多載即可。

但是，如果您希望您的類別在主機/裝置上具有不同的成員，則您將無法在兩個類別中都提供可運作的 H 和 D 多載。在這種情況下，clang 很可能會對您感到不滿。

#ifdef __CUDA_ARCH__
struct S {
  __device__ void foo() { /* use device_only */ }
  int device_only;
};
#else
struct S {
  __host__ void foo() { /* use host_only */ }
  double host_only;
};

__device__ void test() {
  S s;
  // clang generates an error here, because during host compilation, we
  // have ifdef'ed away the __device__ overload of S::foo().  The __device__
  // overload must be present *even during host compilation*.
  S.foo();
}
#endif

我們假設您並不想擁有在 H 和 D 上具有不同成員的類別。例如，如果您要將其中一個類別作為參數傳遞給核心，則它在 H 和 D 上的佈局會有所不同，因此無法正常運作。

為了使此類程式碼與 clang 相容，我們建議您將其分離成兩個類別。如果您需要編寫可在主機和裝置上執行的程式碼，請考慮編寫一個多載的包裝函式，該函式在主機和裝置上返回不同的類型。

struct HostS { ... };
struct DeviceS { ... };

__host__ HostS MakeStruct() { return HostS(); }
__device__ DeviceS MakeStruct() { return DeviceS(); }

// Now host and device code can call MakeStruct().

遺憾的是，這種語法與 nvcc 不相容，因為它不允許您根據 H/D 屬性進行多載。以下是一種適用於 clang 和 nvcc 的語法

struct HostS { ... };
struct DeviceS { ... };

#ifdef __NVCC__
  #ifndef __CUDA_ARCH__
    __host__ HostS MakeStruct() { return HostS(); }
  #else
    __device__ DeviceS MakeStruct() { return DeviceS(); }
  #endif
#else
  __host__ HostS MakeStruct() { return HostS(); }
  __device__ DeviceS MakeStruct() { return DeviceS(); }
#endif

// Now host and device code can call MakeStruct().

希望您不必經常做這種事情。

最佳化 ¶

現代 CPU 和 GPU 在架構上完全不同，因此在 CPU 上運作快速的程式碼在 GPU 上不一定快速。我們對 LLVM 進行了一些變更，使其能夠產生良好的 GPU 程式碼。這些變更包括

直線純量最佳化 – 這些最佳化可減少直線程式碼中的冗餘。
積極的推測執行 – 這主要是為了促進直線純量最佳化，這種最佳化在沿著支配路徑的程式碼上最有效。
記憶體空間推斷 – 在 PTX 中，我們可以操作位於特定「位址空間」（全域、共用、常數或區域）中的指標，或者我們可以操作「通用」位址空間中的指標，它們可以指向任何東西。在非通用位址空間中的操作速度更快，但 CUDA 中的指標沒有明確標註其位址空間，因此由 LLVM 盡可能地推斷它。
繞過 64 位元除法 – 這是我們為 PTX 後端啟用的一種現有最佳化。

在 NVIDIA GPU 上，64 位元整數除法的速度比 32 位元整數除法慢得多。在我們的基準測試中，許多 64 位元除法的除數和被除數在執行時都適合 32 位元。此最佳化為這種常見情況提供了一條快速路徑。
積極的迴圈展開和函數內聯 – 對於 GPU 來說，迴圈展開和函數內聯需要比 CPU 更積極，因為 GPU 中的控制流轉移成本更高。更積極的展開和內聯也促進了其他最佳化，例如常數傳播和 SROA，這些最佳化有時可以將程式碼速度提高 10 倍以上。

（程式設計師可以使用 clang 的迴圈展開編譯指示和 __attribute__((always_inline)) 來強制展開和內聯。）

出版物 ¶

Google 的團隊在 CGO 2016 上發表了一篇論文，詳細介紹了他們對 clang/LLVM 所做的最佳化。請注意，「gpucc」不再是一個有意義的名稱：相關工具現在只是普通的 clang/LLVM。

gpucc：一個開源的 GPGPU 編譯器
Jingyue Wu、Artem Belevich、Eli Bendersky、Mark Heffernan、Chris Leary、Jacques Pienaar、Bjarke Roune、Rob Springer、Xuetian Weng、Robert Hundt
2016 年程式碼生成與最佳化國際研討會 (CGO 2016) 論文集

CGO 演講幻燈片

CGO 教程

取得幫助 ¶

如需有關 LLVM 及其 CUDA 支援的說明，請參閱 LLVM 社群。

文件

參與

其他連結

本頁