clangを使用したCUDAのコンパイル¶

はじめに ¶

このドキュメントでは、clangを使用してCUDAコードをコンパイルする方法について説明し、LLVMとclangのCUDA実装に関する詳細について説明します。

このドキュメントでは、CUDAに関する基本的な知識があることを前提としています。 CUDAプログラミングに関する情報は、CUDAプログラミングガイドにあります。

CUDAコードのコンパイル ¶

前提条件 ¶

CUDAはllvm 3.9以降でサポートされています。 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=…引数を使用してその場所を指定することです。

CUDAコンパイルはLinuxでサポートされています。 MacOSとWindowsでのコンパイルは機能する場合と機能しない場合があり、現在メンテナーはいません。

clangの起動 ¶

CUDAコンパイルのためにclangを起動する方法は、通常の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のコンピュートケ capability。たとえば、コンピュートケ capabilityが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の「間違った側のルール」の解釈により、__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の実装とほぼ互換性がありますが、CUDAコードをclangでコンパイルしていることを具体的に検出したい場合があります。

これは、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 ファイルに対してプリプロセッサを実行し、2つのソースファイルに分割します。ホストのソースコードを含む H と、デバイスのソースコードを含む D です。

コンパイル対象の各GPUアーキテクチャ arch について、以下を実行します。

nvcc properを使用して 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は入力ファイルを複数回解析する必要があるのか疑問に思うかもしれません。なぜ1回だけ解析して、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はHDよりも低い優先度でDを呼び出し、さらに低い優先度でHを呼び出します。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 と互換性を持たせるには、2 つのクラスに分割することをお勧めします。ホストとデバイスの両方で動作するコードを作成する必要がある場合は、ホストとデバイスで異なる型を返すオーバーロードされたラッパー関数の記述を検討してください。

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

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

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

残念ながら、このイディオムは nvcc と互換性がありません。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 ではループ展開と関数インライン化を CPU よりも積極的に行う必要があります。より積極的な展開とインライン化は、定数伝播や 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 コミュニティを参照してください。

ドキュメント

参加する

追加リンク

このページ