使用 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 driver version is insufficient for CUDA runtime version」錯誤。

<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，您可能關心如何讓數值程式碼快速執行。GPU 硬體允許比大多數 CPU 更精細地控制數值運算，但這會導致您需要處理更多的編譯器選項。

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

-ffp-contract={on,off,fast}（在編譯 CUDA 時，主機和裝置上的預設值為 fast）控制編譯器是否發出融合乘加運算。
- off：永不發出 fma 運算，並防止 ptxas 融合乘法和加法指令。
- on：在單一語句中融合乘法和加法，但絕不在語句之間融合（C11 語意）。防止 ptxas 融合其他乘法和加法。
- fast：在任何有利可圖的地方融合乘法和加法，甚至跨語句。不防止 ptxas 融合額外的乘法和加法。
融合乘加指令可能比未融合的等效指令快得多，但由於 fma 中的中間結果未四捨五入，因此此標誌可能會影響數值程式碼。
-fcuda-flush-denormals-to-zero（預設值：off）啟用此選項後，浮點運算可能會將非正規化輸入和/或輸出刷新為 0。對非正規化數字的運算通常比對正規化數字的相同運算慢得多。
-fcuda-approx-transcendentals（預設值：off）啟用此選項後，編譯器可能會發出對更快、近似版本的超越函數的呼叫，而不是使用較慢、完全符合 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 對「wrong-side rule」（請參閱下文）的解讀，它通常在 __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 proper 編譯 D。結果是一個 ptx 檔案，用於 P_arch。

選擇性地，調用 PTX 組譯器 ptxas 以產生一個檔案 S_arch，其中包含用於 arch 的 GPU 機器碼 (SASS)。

調用 fatbin 以將所有 P_arch 和 S_arch 檔案組合到單一「fat binary」檔案 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 檔案組合到單一 fat binary 檔案 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 函數將呼叫 D 函數，其優先順序低於 HD，並且將呼叫 H 函數，其優先順序仍然較低。如果強制呼叫 H 函數，則如果我們為此 HD 函數發出程式碼，則程式格式錯誤。我們將其稱為「wrong-side rule」，請參閱以下範例。

在為主機編譯時，規則是對稱的。

一些範例

__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
}

Wrong-side rule 範例

__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(); }

為了 wrong-side rule 的目的，範本函數的行為也類似於 inline 函數：除非它們被實例化（通常作為調用它們的過程的一部分），否則它們不會被 codegen'ed。

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

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

Clang 針對 Host 和 Device 函數宣告的警告 ¶

當 clang 偵測到使用相同簽名宣告或定義主機 (H) 和裝置 (D) 函數時，它可以發出警告。預設情況下未啟用這些警告。

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

-Wnvcc-compat

在 Host/Device 上使用不同的類別 ¶

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

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

但是，如果您希望您的類別在主機/裝置上具有不同的成員，您將無法在兩個類別中都提供有效的 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 的迴圈展開 pragma 和 __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 社群。

文件

參與貢獻

額外連結

本頁