oneAPI DPC++ コンパイラーとランタイムのアーキテクチャー設計

2022.12.13
その他インテル® oneAPI
DPC++LLVMoneAPISYCL

この記事は、インテル社の許可を得て iSUS (IA Software User Society) が翻訳した、https://github.com/intel/llvm/blob/sycl/sycl/doc/design/CompilerAndRuntimeDesign.md で公開されている 2022年10月14日現在の『oneAPI DPC++ Compiler and Runtime architecture design』の日本語参考訳です。原文は更新される可能性があります。原文と翻訳文の内容が異なる場合は原文を優先してください。

この記事の PDF 版はこちらからご利用になれます。

はじめに

この記事では、DPC++ コンパイラーとランタイム・ライブラリーのアーキテクチャーを説明します。DPC++ 仕様については、こちら (英語) を参照してください。

DPC++ コンパイラー・アーキテクチャー

DPC++ アプリケーションのコンパイルフロー

DPC++ コンパイラーの高レベル・コンポーネントの図
図 1: アプリケーションのビルドフロー

DPC++ コンパイラーは論理的に、ホスト・コンパイラーと多数のデバイス・コンパイラー (サポートされているターゲットごとに 1 つずつ) に分けることができます。Clang ドライバーはコンパイルプロセスを管理し、要求されたターゲットごとにデバイス・コンパイラーを一度呼び出し、次にホスト・コンパイラーを呼び出して SYCL* ソースのホスト領域をコンパイルします。コンパイルとリンクが 1 つのコンパイラー・ドライバーで呼び出される最も単純なケースでは、コンパイルが終了すると、デバイス・オブジェクト・ファイル (実際には LLVM IR ファイル) が llvm-link ツールによりリンクされます。生成される LLVM IR モジュールは、llvm-spirv ツールで SPIR-V* モジュールに変換され、clang-offload-wrapper ツールを使ってホスト・オブジェクト・ファイルにラップされます。すべてのホスト・オブジェクト・ファイルとデバイスコードを含むラップ・オブジェクトの準備ができたら、ドライバーは通常のプラットフォーム・リンカーを呼び出し、「Fat binary (ファットバイナリー)」と呼ばれる最終実行ファイルが生成されます。これは、コマンドラインで指定された各ターゲット用にリンクされたイメージを埋め込んだホスト実行ファイルまたはライブラリーです。

以下の選択に応じて、コンパイルプロセスは異なります。

  • リンクとコンパイルを別々に行う
  • 1 つまたは複数のターゲット用にデバイス SPIR-V* モジュールの事前 (AOT) コンパイルを行う
  • デバイスコードを分割し、1 つのモジュールで実行するのではなく、複数のモジュールに分散させる
  • 静的デバイス・ライブラリーとリンクするこれらのシナリオの詳細は、以降のセクションを参照してください。

SYCL* ソースは、通常の C++ コードとしてコンパイルすることもできます。このモードでは、コードの「デバイス領域」は存在せず、すべてホスト上で実行されます。

デバイス・コンパイラーは、以下のコンポーネントに分けることができます。

  • フロントエンド – 入力ソースを解析し、コードのデバイス領域を「アウトライン化」し、デバイスコードに追加の制限 (例外や仮想呼び出しの禁止など) を適用し、デバイスコードのみの LLVM IR と、カーネル名、パラメーター順、ランタイム・ライブラリーのデータ型などの情報を提供する「Integration header (統合ヘッダー)」を生成します。
  • ミドルエンド – 最初の LLVM IR を変換して、バックエンドで処理されるようにします。現在、ミドルエンドの変換には、2 ～ 3 のパスがあるだけです。
    • オプション: アドレス空間推論パス
    • TBD: ミドルエンドのオプティマイザーは任意の LLVM IR 変換を実行できますが、1 つだけ制限があり、バックエンドのコンパイラーが変換した LLVM IR を処理できる必要があります。
    • オプション: LLVM IR から SPIR-V への変換
  • バックエンド – ネイティブの「デバイス」コードを生成します。図 1 では、「Target-specific LLVM compiler (ターゲット固有の LLVM コンパイラー)」ボックスで示されています。コンパイル時 (AOT コンパイルの場合) または実行時 (JIT コンパイルの場合) に起動されます。

設計上の注意: 現在の設計では、SYCL* デバイス・フロントエンド・コンパイラーを使用して統合ヘッダーを生成していますが、これには 2 つの理由があります。まず、ヘテロジニアスな SYCL* アプリケーションを作成するため、どのようなホスト・コンパイラーでも使用できるようにする必要があります。次に、ホストコンパイルに同じ Clang コンパイラーを使用しても、統合ヘッダーで提供される情報は SYCL* ランタイム実装で使用される（インクルードされる）ので、ホストコンパイルが始まる前にヘッダーが利用可能でなければなりません。

Clang フロントエンドでの SYCL* サポート

Clang フロントエンドでの SYCL* サポートは、次のコンポーネントで構成されます。

  • デバイスコードのアウトライン化。このコンポーネントは、シングルソース中の「デバイスコード」を識別してアウトライン化する役割を担います。
  • SYCL* カーネル関数オブジェクト (ファンクターまたはラムダ) の下位変換。このコンポーネントは、SYCL* カーネル用の OpenCL* カーネル関数インターフェイスを作成します。
  • デバイスコードの診断。このコンポーネントは、デバイスコードの言語制限を実施します。
  • 統合ヘッダーの生成。このコンポーネントは、OpenCL* API を介して SYCL* コードのホスト領域とデバイス領域の結合に必要な情報を出力します。

デバイスコードのアウトライン化

以下は、コンパイラーのアウトライン化を示す SYCL* プログラムのコード例です。

int foo(int x) { return ++x; }
int bar(int x) { throw std::exception{"CPU code only!"}; }
...
using namespace sycl;
queue Q;
buffer<int, 1> a{range<1>{1024}};
Q.submit([&](handler& cgh) {
      auto A = a.get_access<access::mode::write>(cgh);
      cgh.parallel_for<init_a>(range<1>{1024}, [=](id<1> index) {
        A[index] = index[0] * 2 + foo(42);
      });
    }
...

この例では、コンパイラーは sycl::handler::parallel_for メソッドに渡されるラムダ式と、デバイス用のラムダ式から呼び出される関数 foo をコンパイルする必要があります。

また、コンパイラーは単一のソースコードの「デバイス」領域をコンパイルする際に、ソースコードのデバイス領域 (sycl::handler::parallel_for に渡されるラムダ式の内容や、このラムダ式から呼ばれる関数) で使われていない bar 関数を無視する必要があります。

現在のアプローチは、ランタイムの SYCL* カーネル属性を使用して、sycl::handler::parallel_for に渡されるコードを「カーネル関数」としてマークします。ランタイム・ライブラリーは foo を「デバイス」コードとしてマークすることはできません。コンパイラーが、カーネル関数からアクセスできるすべてのシンボルをトラバースし、それらを新しい SYCL* デバイス属性でマークしてコードの「デバイス領域」に追加する必要があります。

ラムダ関数オブジェクトと名前付き関数オブジェクトの下位変換

ホストとデバイス間で共有されるすべての SYCL* メモリー・オブジェクト (バッファー/イメージ、これらのオブジェクトは OpenCL* バッファーとイメージにマッピングされる) は、特別な accessor クラスを介してアクセスする必要があります。これらのクラスの「デバイス」側の実装には、デバイスメモリーへのポインターが含まれます。OpenCL* にはカーネル引数として内部にポインターを持つ構造体を渡す方法がないため、ホストとデバイスの間で共有されるすべてのメモリー・オブジェクトは、生のポインターとしてカーネルに渡されなければなりません。

SYCL* はまた、ホストからデバイスにカーネル引数を渡す特別なメカニズムを持っています。OpenCL* では、カーネル引数ごとに clSetKernelArg 関数を呼び出してカーネル引数を設定します。一方、SYCL* では、すべてのカーネル引数は「SYCL* カーネル関数」のフィールドで、ラムダ関数または名前付き関数オブジェクトとして定義でき、カーネルを呼び出す SYCL* 関数 (parallel_for や single_task など) への引数として渡されます。例えば、上記のコード例では accessor A がキャプチャーされたカーネル引数です。

SYCL* カーネル・データ・メンバーの OpenCL* カーネル引数へのマッピングを容易にし、OpenCL* の制限に対応するため、コンパイラー内部に OpenCL* カーネル関数を生成する機能を追加しました。OpenCL* カーネル関数は、SYCL* カーネル関数の本体を含み、OpenCL* 形式のパラメーターを受け取り、さらに、これらのパラメーターで SYCL* カーネル・データ・メンバーを初期化する操作を行います。以下は、上記のコード例の OpenCL* カーネル関数の擬似コードです。

// SYCL* カーネルは SYCL* ヘッダーで定義されている:
template <typename KernelName, typename KernelType/*, ...*/>
__attribute__((sycl_kernel)) void sycl_kernel_function(KernelType KernelFuncObj) {
  // ...
  KernelFuncObj();
}

// OpenCL* カーネル関数を生成する
__kernel KernelName(global int* a) {
  KernelType KernelFuncObj; // カーネル関数オブジェクトの宣言は AST では名前がない
  // カーネル関数オブジェクトのアクセサー A という
  // キャプチャー・フィールドを引数の
  // グローバルポインターで初期化する
  KernelFuncObj.A.__init(a);
  // SYCL* ヘッダーからの SYCL* カーネル本体
  {
    KernelFuncObj();
  }
}

OpenCL* カーネル関数は、AST (抽象構文木) ノードを使用して Sema 内のコンパイラーによって生成されます。カーネル・パラメーターの受け渡しの詳細は、「SYCL* カーネル・パラメーターの処理と配列のサポート」 (英語) を参照してください。

ドライバーの SYCL* サポート

ドライバーの SYCL* オフロードサポートは、Clang ドライバーの概念に基づいており、以下を定義します。

  • ターゲットトリプルと各ターゲットのネイティブ・ツールチェーン (SPIR-V* のような「仮想」ターゲットを含む)
  • 一般的なオフロードアクションに基づく SYCL* オフロードアクション

SYCL* コンパイル・パイプラインには、ほかのコンパイルシナリオと比較して特異な点があります。パイプライン内の一部のアクションは、ファイルの複数の「クラスター」を出力し、後でほかのアクションによって処理される場合があります。例えば、各デバイスバイナリーには、シンボルテーブルと特殊化定数マップ (SYCL* ランタイム・ライブラリーによって使用される追加情報) が含まれる場合があり、オフロード・ラッパー・ツールでデバイスバイナリー記述子に格納する必要があります。デバイスコード分割機能を有効にすると、このようなファイルのセット (クラスター) が複数出力される可能性があります (デバイスバイナリーごとに 1 つ)。

現在の Clang ドライバーの設計では、以下をモデル化できません。

  1. アクショングラフにおける複数の入力/出力
  2. 複数の入力/出力の論理的なグループ化。例えば、入力や出力は複数のファイルのペアで構成され、それぞれのペアは 1 つのデバイス・コード・モジュールの情報を表します [デバイスコードのファイル、エクスポートされたシンボルのファイル]

これをサポートするため、SYCL* は file-table-tform ツールを導入しています。このツールは、入力引数として渡されたコマンドに従ってファイルテーブルを変換します。テーブルの各行はファイルのクラスターを表し、各列はクラスターに関連するデータの種類を表します。このツールはカラムを置き換えたり、抽出することができます。例えば、sycl-post-link ツールは、2 つのファイルクラスターと、そのクラスター内のすべてのファイルを参照する以下のようなファイルテーブルを出力できます。

  [Code|Symbols|Properties]
  a_0.bc|a_0.sym|a_0.props
  a_1.bc|a_1.sym|a_1.props

アクショングラフでは、このツールはファイルテーブル (TY_Tempfiletable Clang 入力タイプ) やファイルリスト (TY_Tempfilelist) を入力として使用し、要求された変換を行い、ファイルテーブルやリストを出力します。実際には入力と出力は複数ですが、Clang の設計上は 1 つのままです。

例えば、コンパイルオプションによっては、上記の「Code」列のファイルは、sycl-post-link ツールのコード変換シーケンスの一部として行われる、デバイスコードの分割ステップの後に AOT コンパイルを行う必要があるかもしれません。その場合、ドライバーは以下の処理を行います。

  • file-table-tform を使って、コードファイルを抽出し、ファイルリストを作成します。 
    a_0.bc
a_1.bc
  • このファイルリストを AOT コンパイルコマンドと一緒に llvm-foreach ツールに渡して、リストにあるすべてのファイルに対してコンパイルコマンドを実行します。これにより、別のファイルリストが生成されます。 
    a_0.bin
a_1.bin
  • 再度 file-table-tform を呼び出して、filetable の .bc を .bin に置き換えて、新しい filetable を取得します。 
      [Code|Symbols|Properties]
  a_0.bin|a_0.sym|a_0.props
  a_1.bin|a_1.sym|a_1.props
  • 最後に、この filetable を clang-offload-wrapper ツールに渡して、これらのファイルすべてを埋め込んだラッパー・オブジェクトをビルドします。

テーブルに行 (クラスター) や列 (「マニフェスト」ファイルなど) が追加されても、グラフは変わりません。

SYCL* オフロードの有効化

SYCL* 仕様に記載されている SMCP (Single-Source Multiple Compiler-Passes) 技術に従ってコンパイルするには、Clang ドライバーに特別なオプションを渡す必要があります。

-fsycl

このオプションを指定すると、ドライバーはホスト・コンパイラーと -fsycl-targets オプションで指定されたターゲット用のいくつかの SYCL* デバイス・コンパイラーを起動します。-fsycl-targets を指定しないと、単一の SPIR-V* ターゲットが想定され、このターゲット用の単一のデバイス・コンパイラーが起動されます。

-sycl-std オプションでコンパイルに使用する SYCL* 標準のバージョンを指定できます。デフォルト値は 1.2.1 です。

AOT コンパイル

事前 (Ahead of Time、略称 AOT) コンパイルは、最終コード生成をアプリケーションの実行時まで延期する JIT コンパイルとは対照的に、コンパイル時にバックエンドを呼び出して最終バイナリーを生成します。

AOT コンパイルは、JIT コンパイルをスキップすることでアプリケーションの実行時間を短縮し、最終的なデバイスコードを展開前にテストできるようにします。

JIT コンパイルは、デバイスコードの移植性とターゲット固有の最適化を提供します。

ネイティブターゲットのリスト

AOT コンパイルモードでは、デバイスコードをコンパイルするターゲット・アーキテクチャーのセットを指定する必要があります。デフォルトでは、コンパイラーは SPIR-V* と OpenCL* デバイスの JIT コンパイラーを生成し、ネイティブ・ターゲット・バイナリーを生成します。

ターゲットトリプル triple1 と triple2 で識別されるターゲット・アーキテクチャー用のバイナリーを生成するには、次の SYCL* コンパイラー・オプションを使用します。

-fsycl-targets=triple1,triple2

SYCL* カーネルから、2 つのターゲットトリプルで識別されるデバイス用のバイナリーが生成されます。このオプションは、SYCL* カーネルコードをコンパイルするため、どのデバイス・コンパイラーを呼び出す必要があるかをドライバーに指示します。デフォルトでは、JIT コンパイルが想定されており、デバイスコードは 1 つのターゲットトリプル ([spir,spir64]-*-*) 向けにコンパイルされます。

デバイスコード形式

各デバイス・コンパイラーは、いくつかのコード形式をサポート可能で、特定のコード形式を指定するニーモニックを定義して解釈することができます。例えば、「visa:3.3」は、インテル® GPUターゲット (Gen アーキテクチャー) 向け仮想 ISA バージョン 3.3 とすることができます。ユーザーは、OpenMP* と同様に、ターゲット固有のオプション構文を使用してコード形式を指定できます。

-Xsycl-target-backend=<triple> "arg1 arg2 ..."

例えば、Gen9/vISA3.3 へのオフロードをサポートする場合、以下のオプションを使用します。

-fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device skl"

ドライバーは -device skl パラメーターを解析することなく、直接 Gen デバイスのバックエンド・コンパイラー ocloc に渡します。

ocloc は、いくつかの ISA バージョン/Gen アーキテクチャー向けのオフラインコンパイルにも対応しています。例えば、すべての第 9 世代インテル® GPU (Gen9) プラットフォームと互換性のあるデバイスバイナリーを作成するには、以下のオプションを使用します。

-fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device gen9"

サポートされているプラットフォームと引数の構文の詳細は、ローカルの ocloc インストールを検出して ocloc compile --help を実行し、GPU オフライン・コンパイラーのマニュアルを参照してください。

リンクとコンパイルを別々に行う

コンパイラーは以下をサポートしています。

  • バックエンドに供給する最終的な SPIR-V* を生成する前に、異なるソースファイルから取得したデバイスコードをリンクする
  • アプリケーションのビルドをコンパイルステップとリンクステップに分離する

全体的なビルドフローは、図 1 に示したものと比較して、次のように変化します。コンパイルステップでは、オフロードバンドラーを使用して、同じヘテロジニアス・ソースから生成された各<ホスト・オブジェクト, デバイスコード IR> ペアに対して、「ファット・オブジェクト」を生成します。コンパイル結果のファット・オブジェクト・ファイルは、通常の非オフロード・コンパイラーによって生成されるオブジェクト・ファイルに似ています。リンクステップでは、入力されたファット・オブジェクトを構成要素に分解し、図 1 と同様に、ホストコードとデバイスコードを別々にリンクして最終的に「ファットバイナリー」を生成します。

下の図は、ビルドフローの違いを示しています。オフロードバンドラー/アンバンドラーのアクションは、基本的に図 1 に示すように llvm-link と linker 呼び出しの間に挿入されます。

複数のソースのコンパイルフロー
図 2. コンパイルとリンクの分離

現在の実装では、LLVM IR を fat objects のデフォルトのデバイスバイナリー形式として使用し、「リンクされた LLVM IR」を SPIR-V* に変換しています。この理由の 1 つは、SPIR-V* がテンプレート関数のリンクに対応していないため、複数のモジュールで定義されている可能性があり、リンカーが複数の定義を解決しなければならないからです。LLVM IR は、SPIR-V* にはない「1 定義ルール」を満たすため、関数属性を使用します。

ファットバイナリー作成の詳細

「ファットバイナリー」とは、ホストリンクの最終段階で生成される、デバイスバイナリーを埋め込んだホストバイナリーです。実行すると、SYCL* ランタイム・ライブラリー内で利用可能なすべてのデバイスバイナリーが自動的に登録されます。このセクションでは、この手順を説明します。

出力されるファットバイナリーは、通常のリンカー (例えば、Linux* では ld、Windows* では link.exe) で作成されます。リンカーがデバイスバイナリーを埋め込むことができるように、まず「ラッパー・オブジェクト」と呼ばれるホスト・オブジェクト・ファイルに「ラップ」します。ラッパー・オブジェクトは、ほかのホスト・オブジェクトやライブラリーと通常どおりにリンクされます。

ラッパー・オブジェクトは、clang-offload-wrapper ツール (オフロードラッパー) によって作成されます。作成されたラッパー・オブジェクトには、2 つの主要な構成要素があります。

  1. オブジェクトのデータセクションにある特別なデータ構造を指すグローバルシンボル (オフロード記述子)。バイナリー、バイナリーの数、各バイナリーが定義するシンボルなど、ラップされたデバイスバイナリーに関するすべての必要な情報が含まれます。
  2. 登録/解除関数。最初の関数は、実行時に親ファットバイナリーがプロセスにロードされたときに呼び出されるように特別なセクションに配置され、2 つ目の関数は、親ファットバイナリーがアンロードされたときに呼び出されるように別のセクションに配置されます。登録関数は基本的にオフロード記述子へのポインターを受け取り、それをパラメーターとして SYCL* ランタイム・ライブラリーの登録関数を呼び出します。

オフロード記述子の型階層は pi.h ヘッダーに記述されています。トップレベルの構造体は pi_device_binaries_struct です。

デバイスリンク

-fsycl-link コンパイラー・オプションは、ホストコードを完全にリンクせずにデバイスコードを完全にリンクするようにコンパイラーに指示します。コンパイル結果は、完全にリンクされたデバイスバイナリーを含むファット・オブジェクトになります。これにより、ホストコードにのみ影響する変更を行う際に、再コンパイル時間を節約できます。デバイスイメージの生成に長い時間がかかる場合 (FPGA など)、これは大いに役立ちます。

例えば、ソースコードを dev_a.cpp、dev_b.cpp、host_a.cpp、host_b.cpp の 4 ファイルに分割し、dev_a.cpp と dev_b.cpp にのみデバイスコードが含まれる場合、コンパイル処理を 3 ステップに分割することが可能です。

  1. デバイスリンク: dev_a.cpp dev_b.cpp -> dev_image.o (デバイスイメージを含む)
  2. ホストコンパイル (c): host_a.cpp -> host_a.o; host_b.cpp -> host_b.o
  3. リンク: dev_image.o host_a.o host_b.o -> 実行ファイル

ステップ 1 は、ターゲットによっては数時間かかる場合があります。host_a.cpp と host_b.cpp だけを変更して再コンパイルする場合、時間のかかるステップ 1 はスキップして、ステップ 2 と 3 だけを実行できます。

デバイス・リンク・ステップに必要なすべての関連ファイルが揃っていることを確認するのは、コンパイラーの責任です。以下の 2 つのケースをチェックする必要があります。

  1. デバイス・リンク・ステップに存在するカーネルによって参照されるシンボルがない場合 (例えば、既知のカーネルによって呼び出される関数や使用されるグローバル変数など)
  2. カーネルがない場合

ケース 1 は、デバイスバイナリー生成段階 (ステップ 1) で既知のカーネルをスキャンすることで確認できます。ケース 2 は、最後のリンク段階 (ステップ 3) でドライバーが新たに導入されたカーネルをチェックして検証する必要があります。

llvm-no-spir-kernel ツールは、ドライバーによるケース 2 のチェックを容易にするために導入されました。これは、モジュールにカーネルが含まれるかどうかを検出するもので、次のように起動します。

llvm-no-spir-kernel host.bc

カーネルが存在しない場合は 0 を、それ以外の場合は 1 を返します。

デバイス・コード・リンク後のステップ

リンク時に、すべてのデバイスコードは常に単一の LLVM IR モジュールにリンクされます。sycl-post-link ツールは、オフロードラッパーに渡す前に、この LLVM IR モジュール上で次のようないくつかの最終変換を実行します。

  • デバイスコードの分割
  • シンボルテーブルの生成
  • 特殊化定数の下位変換

オプションに応じて、sycl-post-link は、単一の LLVM IR ファイル、または、複数のファイルとそれらすべてを参照するファイルテーブルを出力します。ファイルテーブルについては、「ドライバーの SYCL* サポート」セクションを参照してください。下の図は、単一のリンクされた LLVM IR モジュールからラッパー・オブジェクトを作成するため、コンパイルプロセスがたどる可能性のある Clang アクショングラフを示したものです。以下の要因に応じて、このグラフには複数のバリエーションがあります。

  • 特定のターゲット要件
  • デバイスコードの分割
  • AOT コンパイル

複数のソースのコンパイルフロー
図 3: デバイスコードのリンクフロー

グラフのエッジの色は、上記の要因によってどのような経路をたどるかを示しています。また、各エッジには入出力ファイルの種類が注釈されています。この図では、分かりやすくするため llvm-foreach ツールの起動を表示していません。このツールは、ファイルリスト中の各ファイルに対して与えられたコマンドラインを呼び出します。この図では、入出力タイプが TY_tempfilelist でターゲットが PTX でない場合、llvm-spirv と AOT バックエンドにこのツールが適用されます。続く、file-table-tform は 2 つの入力 (ファイルテーブルと llvm-spirv または AOT バックエンドからのファイルリスト) を受け取ります。PTX ターゲット処理は現在、単一の入力ファイルしか受け付けていないので、file-table-tform でファイルテーブルからコードファイルを抽出して、「PTX ターゲット処理」ステップで処理します。生成されたデバイスバイナリーは、file-table-tform で抽出したコードファイルの代わりにファイルテーブルに挿入されます。

デバイスコードの分割

以下の場合、すべてのデバイスコードを単一の SPIR-V* モジュールに含めないほうがよいでしょう。

  1. 何千ものカーネルが定義されており、そのうちのごく一部のみが実行時に使用される。1 つの SPIR-V* モジュールにすべてを含めると、JIT 時間が大幅に増加します。
  2. デバイスコードが異なるデバイス向けに特殊化される。例えば、FPGA で実行するカーネルのみが FPGA でのみ利用可能な拡張を使用できるため、ほかのデバイスでは、この特定のカーネルが呼び出されなくても、JIT コンパイルに失敗します。

この問題を解決するため、コンパイラーは 1 つのモジュールをより小さなモジュールに分割できます。以下がサポートされています。

  • ソース (変換ユニット) ごとに個別のモジュールを生成する
  • カーネルごとに個別のモジュールを生成する

現在のアプローチは以下のとおりです。

  • SYCL* フロントエンドで、カーネルごとに変換ユニット ID を含む特別なメタデータを生成します。この ID は、変換ユニットごとにカーネルをグループ化するのに使用されます。
  • llvm-link を使用して、すべてのデバイスの LLVM モジュールをリンクします。
  • 完全にリンクされたモジュールを分割します。
  • 実行時に適切なモジュールが選択されるように、生成されるデバイスモジュールごとにシンボルテーブル (カーネルのリスト) を生成します。
  • 生成されるモジュールごとに SPIR-V* 変換と AOT コンパイル (要求された場合) を実行します
  • 提供されたカーネルに関する情報を、各デバイスイメージのラップ・オブジェクトに追加します。

デバイスコードの分割

デバイスコードの分割プロセス

「split (分割)」ボックスは、専用ツール sycl-post-link の機能として実装されています。このツールは、入力モジュールを分割するため一連の LLVM パスを実行し、生成されるデバイスモジュールごとにシンボルテーブル (カーネルのリスト) を生成します。

デバイスコードの分割を有効にするには、Clang ドライバーに次のオプションを渡す必要があります。

-fsycl-device-code-split=<value>

value に設定可能な値は 3 つあります。

  • per_source – ソース (変換ユニット) ごとに個別のモジュールを生成します。
  • per_kernel – カーネルごとに個別のモジュールを生成します。
  • off – デバイスコードの分割を無効にします。
シンボルテーブルの生成

TBD

特殊化定数の下位変換

https://github.com/intel/llvm/blob/sycl/sycl/doc/design/SpecializationConstants.md (英語) を参照してください。

CUDA* サポート

nvptx64-nvidia-cuda が -fsycl-targets に渡された場合、ドライバーは NVPTX へのコンパイルをサポートします。

ほかの AOT ターゲットとは異なり、中間コンパイル・オブジェクトからリンクされたビットコード・モジュールが SPIR-V* ステップを通過することはありません。代わりに、ビットコード形式で直接 NVPTX バックエンドに渡されます。すべての生成されたビットコードは、libdevice.bc (CUDA* SDKによって提供) と libspirv-nvptx64--nvidiacl.bc バリアント (libclc プロジェクトによってビルド) の 2 つのライブラリーに依存します。libspirv-nvptx64--nvidiacl.bc は直接使用されません。代わりに、Linux* と Windows* の間のプリミティブ型の違いを処理する、リマングルバリアント remangled-l64-signed_char.libspirv-nvptx64--nvidiacl.bc と remangled-l32-signed_char.libspirv-nvptx64--nvidiacl.bc を生成するために使用されます。

CUDA* 向けのデバイス・コード・リンク後のステップ

「デバイス・コード・リンク後のステップ」の PTX ターゲット処理中に、CUDA* ターゲット用の llvm ビットコード・オブジェクトが共通の llvm-link ステップでリンクされ、その後 sycl-post-link ツールを使用して分割されます。一時的なビットコード・ファイルごとに Clang が呼び出され、libspirv-nvptx64--nvidiacl.bc と libdevice.bc をリンクし、生成されたモジュールを NVPTX バックエンドを使用して PTX にコンパイルします。このPTX ファイルから ptxas ツール (CUDA* SDK の一部) を使用して cubin を生成し、PTX ファイルと cubin から fatbinary を使用して CUDA* fatbin を生成します。CUDA* fatbin は、sycl-post-link によって生成されたファイルテーブル内の llvm ビットコード・ファイルを置き換え、テーブルはオフロード・ラッパー・ツールに渡されます。

NVPTX AOT ビルド

コンパイラーが NVPTX をターゲットにしているかどうかのチェック

SYCL* コンパイラーがデバイスモードで NVPTX バックエンドをターゲットにしている場合、コンパイラーは __SYCL_DEVICE_ONLY__ マクロと __NVPTX__ マクロを定義しています。このマクロの組み合わせは、SYCL* カーネルで NVPTX 固有のコードパスを有効にするために安全に使用することができます。

注: これらのマクロは、デバイスコンパイル時にのみ定義されます。

NVPTX ビルトイン

ビルトインは OpenCL* C の libclc 内に実装されています。OpenCL* C は long 型を 64 ビットとして扱い、long long 型を持たないのに対し、Windows* DPC++ は long 型を 32 ビット整数のように扱い、long long 型を 64 ビット整数のように扱います。プリミティブ型の違いは、アプリケーションが互換性のない libclc ビルトインを使用する原因になる場合があります。リマングラーは、Windows* と Linux* の両方をサポートするため、異なるリマングル関数名を持つ複数の libspriv ファイルを作成します。CUDA* バックエンドをターゲットとする SYCL* アプリケーションをビルドする場合、ドライバーは、ホストターゲットが Windows* の場合は remangled-l32-signed_char.libspirv-nvptx64--nvidiacl.bc でデバイスコードをリンクし、ホストターゲットが Linux* の場合は remangled-l64-signed_char.libspirv-nvptx64--nvidiacl.bc でデバイスコードをリンクします。

SYCL* コンパイラーがデバイスモードで、NVPTX バックエンドをターゲットにしている場合、コンパイラーは Clang でサポートされている NVPTX ビルトインを公開します。

注: これは、ほかのターゲットやホストではサポートされていない NVPTX 固有の機能を有効にします。

例:

double my_min(double x, double y) {
#if defined(__NVPTX__) &amp;&amp; defined(__SYCL_DEVICE_ONLY__)
  // デバイスモードで、NVPTX ターゲット向けに
  // コンパイルする場合のみ利用可能
  return __nvvm_fmin_d(x, y);
#else
  return x &lt; y ? x : y;
#endif
}
ローカル・メモリー・サポート

CUDA* では、ユーザーはホストに割り当てられた共有メモリー (SYCL* のローカルアクセサーにマッピングされる) の 1 チャンクのみを割り当てることができます。このメモリーチャンクは、配列 extern __shared__ <type> <name>[]; として割り当てられ、LLVM では CUDA* 共有メモリーアドレス空間への外部グローバルシンボルとして表現されます。NVPTX バックエンドは、これを .extern .shared .align 4 .b8 PTX 命令に下位変換します。

SYCL* では、ユーザーは複数のローカルアクセサーを割り当てて、それらをカーネル・パラメーターとして渡すことができます。SYCL* フロントエンドは SYCL* カーネル呼び出しを OpenCL* 準拠のカーネルエントリーに変換する際に、ローカルアクセサーを OpenCL* ローカルメモリー (CUDA* 共有メモリー) へのポインターに変換しますが、これは CUDA* カーネルでは無効です。

CUDA* 向けの SYCL* 変換を有効なものにするため、CUDA* 固有パスで SYCL* は次のことを行います。

  • CUDA* 共有メモリーアドレス空間へのグローバルシンボルを作成します。
  • CUDA* 共有メモリーへのすべてのポインターを、グローバルシンボルで使用するオフセットをバイト単位で表す 32 ビット整数に変換します。
  • 変換されたポインターのすべての使用を、パラメーターとして渡された整数値でオフセットされたグローバルシンボルへのアドレスに置き換えます。

例えば、次のカーネルについて考えてみます。

define void @SYCL_generated_kernel(i64 addrspace(3)* nocapture %local_ptr, i32 %arg, i64 addrspace(3)* nocapture %local_ptr2) {
  %0 = load i64, i64 addrspace(3)* %local_ptr
  %1 = load i64, i64 addrspace(3)* %local_ptr2
}

CUDA* をターゲットとする場合、上記のコードは次のように変換されます。

@SYCL_generated_kernel.shared_mem = external dso_local local_unnamed_addr addrspace(3) global [0 x i8], align 4

define void @SYCL_generated_kernel(i32 %local_ptr_offset, i32 %arg, i32 %local_ptr_offset2) {
  %new_local_ptr = getelementptr inbounds [0 x i8], [0 x i8] addrspace(3)* @SYCL_generated_kernel.shared_mem, i32 0, i32 %local_ptr_offset
  %new_local_ptr2 = getelementptr inbounds [0 x i8], [0 x i8] addrspace(3)* @SYCL_generated_kernel.shared_mem, i32 0, i32 %local_ptr_offset2
  %0 = load i32, i32 addrspace(3)* %new_local_ptr
  %1 = load i64, i64 addrspace(3)* %new_local_ptr2
}

ランタイム側では、ローカルメモリー引数を設定する場合、CUDA* API 実装は、使用されたローカルメモリーの累積サイズに対するオフセットとして引数を内部的に設定します。この方法は、既存の API インターフェイスを維持します。

グローバル・オフセット・サポート

CUDA* API は、SYCL* が期待するグローバル・オフセット・パラメーターをネイティブでサポートしていません。

これをエミュレートし、生成されるカーネルを準拠させるため、NVPTX バックエンドでこの暗黙のパラメーターの使用を具体化する llvm.nvvm.implicit.offset 組込み関数 (Clang ビルトインの __builtin_ptx_implicit_offset) が導入されました。AMDGCN は llvm.amdgpu.implicit.offset や __builtin_amdgcn_implicit_offset と同じアプローチを使用します。この組込み関数は、3 つの要素の配列を参照する i32 へのポインターを返します。

コールグラフで暗黙のオフセット組込み関数に到達する非カーネル関数は、i32 へのポインター型の暗黙のパラメーターで拡張されます。この組込み関数を使用してこれらの関数を呼び出すカーネルはクローンされます。

  • オリジナルカーネルは 3 つの i32 配列を 0 に初期化し、この配列へのポインターを暗黙のパラメーターとして各関数に渡します。
  • クローンされた関数型は、3 つの i32 配列型の暗黙のパラメーターで拡張されます。この配列へのポインターは、暗黙のパラメーターを持つ各関数に渡されます。

ランタイムは両方のカーネルを照会し、以下のロジックに基づいて適切なものを呼び出します。

  • 2 つのバージョンが存在する場合、グローバルオフセットが 0 であればオリジナルのカーネルが呼び出され、そうでなければクローンされたものが呼び出され、オフセットを値 (CUDA* バックエンドの場合)、または参照 (AMD の場合) で渡します。
  • 関数が 1 つしか存在しない場合、カーネルはこのパラメーターを使用しないため、無視されます。

例えば、次のコードについて考えてみます。

declare i32* @llvm.nvvm.implicit.offset()

define weak_odr dso_local i64 @other_function() {
  %1 = tail call i32* @llvm.nvvm.implicit.offset()
  %2 = getelementptr inbounds i32, i32* %1, i64 2
  %3 = load i32, i32* %2, align 4
  %4 = zext i32 %3 to i64
  ret i64 %4
}

define weak_odr dso_local void @other_function2() {
  ret
}

define weak_odr dso_local void @example_kernel() {
entry:
  %0 = call i64 @other_function()
  call void @other_function2()
  ret void
}

上記のコードは、次のように変換されます。

define weak_odr dso_local i64 @other_function(i32* %0) {
  %2 = getelementptr inbounds i32, i32* %0, i64 2
  %3 = load i32, i32* %2, align 4
  %4 = zext i32 %3 to i64

  ret i64 %4
}

define weak_odr dso_local void @example_kernel() {
entry:
  %0 = alloca [3 x i32], align 4
  %1 = bitcast [3 x i32]* %0 to i8*
  call void @llvm.memset.p0i8.i64(i8* nonnull align 4 dereferenceable(12) %1, i8 0, i64 12, i1 false)
  %2 = getelementptr inbounds [3 x i32], [3 x i32]* %0, i32 0, i32 0
  %3 = call i64 @other_function(i32* %2)
  call void @other_function2()
  ret void
}

define weak_odr dso_local void @example_kernel_with_offset([3 x i32]* byval([3 x i32]) %0) {
entry:
  %1 = bitcast [3 x i32]* %0 to i32*
  %2 = call i64 @other_function(i32* %1)
  call void @other_function2()
  ret void
}

注: 現時点で、カーネルの命名規則は完全ではありません。

SPIR-V* 形式との統合

このセクションでは、C++ のクラスと関数から SPIR-V* 固有の型と操作を生成する方法について説明します。

SYCL* C++ プログラムからヘテロジニアス・システムで実行可能なコードへの変換は、3 つのステップで構成されていると考えることができます。

  1. SYCL* C++ プログラムの LLVM IR への変換
  2. LLVM IR から SPIR-V* への変換
  3. SPIR-V* からマシンコードへの変換

LLVM-IR から SPIR-V* への変換は、専用ツール translator (英語) で行われます。このツールは、通常の LLVM IR の型や操作などの大部分を SPIR-V* に正しく変換します。

例:

  • 型: i32 → OpTypeInt
  • 操作: load → OpLoad
  • 呼び出し: call → OpFunctionCall

SPIR-V* は、LLVM IR に該当するものがない特殊な組込み型や操作を定義します。以下に例を示します。

  • 型: ??? → OpTypeEvent
  • 操作: ??? → OpGroupAsyncCopy

LLVM IR から SPIR-V* への特殊型の変換もサポートされていますが、いくつかの特別な要件に準拠している必要があります。残念ながら、LLVM IR には、特殊な組込み型と操作の正規形がなく、また、OpenCL* C フロントエンド・コンパイラーが生成した既存の表現を再利用することもできません。例えば、OpenCL* C フロントエンド・コンパイラーが生成した LLVM IR で、OpGroupAsyncCopy 操作は次のようになります。

@_Z21async_work_group_copyPU3AS3fPU3AS1Kfjj(float addrspace(3)*, float addrspace(1)*, i32, i32)

これは通常の関数であり、C++ ソースから生成されたユーザーコードと競合する可能性があります。

DPC++ コンパイラーは、OpenCL* C++ コンパイラーのプロトタイプ向けに開発された変更されたソリューションを使用しています。

ソリューションは、サンプラー、イベント、イメージタイプなどの OpenCL* データ型を再利用しますが、C++ コードとの潜在的な競合を避けるため異なるスペル (つづり) を使用します。SYCL* モードで有効な OpenCL* 型のスペル規則は以下のとおりです。

__ocl_&lt;OpenCL_type_name&gt; // 例: __ocl_sampler_t、__ocl_event_t

OpenCL* 型を使用する操作は、こちらのドキュメント (英語) で説明されている特別な命名規則を使用します。これより、SPIR-V* 変換器での SYCL* の特殊化を避け、OpenCL* 型向けに開発された Clang インフラストラクチャーを活用できます。

LLVM に該当するものがない SPIR-V* 操作は、ヘッダーで宣言され (定義はされない)、以下の要件を満たします。

  • 操作は C++ で表現され、C++ の例外を発生させない extern 関数であること
  • C++ で定義されていないこと

例えば、以下の C++ コードは、SPIR-V* の OpGroupAsyncCopy という操作として認識され、変換されます。

template &lt;typename dataT&gt;
extern __ocl_event_t
__spirv_OpGroupAsyncCopy(int32_t Scope, __local dataT *Dest,
                         __global dataT *Src, size_t NumElements,
                         size_t Stride, __ocl_event_t E) noexcept;

__ocl_event_t e =
  __spirv_OpGroupAsyncCopy(cl::__spirv::Scope::Workgroup,
                           dst, src, numElements, 1, E);

SPIR-V* の特殊な型と操作を使用する際の詳細と規則

SPIR-V* 固有の C++ 列挙子とクラスは、sycl/include/CL/__spirv/spirv_types.hpp ファイルで宣言されます。

SPIR-V* 固有の C++ 関数の宣言は、sycl/include/CL/__spirv/spirv_ops.hpp ファイルにあります。

SPIR-V* 固有の関数は、SYCL* ホストデバイス向けに sycl/source/spirv_ops.cpp ファイルで実装されています。

アドレス空間の扱い

SYCL* 仕様では、標準 C++ ツールチェーンと SYCL* コンパイラー・ツールチェーンでコンパイルできるように、アクセラレーター上の非連続メモリー領域へのポインターを C++ クラスで表現しています。

例:

// SYCL* モードがオンで、非標準の修飾が使用可能なことを確認する
#if defined(__SYCL_DEVICE_ONLY__)
// GPU/アクセラレーターの実装
template &lt;typename T, address_space AS&gt; class multi_ptr {
  // DecoratedType は対応するアドレス空間属性を型 T に適用する
  // DecoratedType&lt;T, global_space&gt;::type == &quot;__attribute__((opencl_global)) T&quot;
  // 詳細は sycl/include/sycl/access/access.hpp を参照
  using pointer_t = typename DecoratedType&lt;T, AS&gt;::type *;
 
  pointer_t m_Pointer;
  public:
  pointer_t get() { return m_Pointer; }
  T&amp; operator* () { return *reinterpret_cast&lt;T*&gt;(m_Pointer); }
}
#else
// CPU/ホスト実装
template &lt;typename T, address_space AS&gt; class multi_ptr {
  T *m_Pointer; // 通常の非修飾ポインター
  public:
  T *get() { return m_Pointer; }
  T&amp; operator* () { return *m_Pointer; }
}
#endif

コンパイラーのモードに応じて、multi_ptr は内部データをアドレス空間属性で装飾するかしないかを決定します。

SYCL* モードのメインアドレス空間のセマンティクスは、明示的なアドレス空間属性なしで宣言の型に OpenCL* 汎用アドレス空間を割り当てないという点で、OpenCL* と異なります。OpenMP*/CUDA*/HIP など、ほかのシングルソースの C++ ベースの GPU プログラミング・モードと同様に、SYCL* はアドレス空間属性のない型に Clang の「デフォルト」のアドレス空間を使用します。LLVM IR に移行する際、デフォルトのアドレス空間は SPIR* の一般的なアドレス空間にマッピングされます。宣言は宣言内容に応じて、関連するメモリー領域に割り当てられ、それらへのポインターは generic にキャストされます。この設計には、C++ との型システムの一貫性を保つ一方、ツールに SPIR* メモリーモデル (およびほかの GPU ターゲット) と整合性のあるデバイスコードの生成を可能にする、という 2 つの重要な特徴があります。

関数内部の変数宣言について考えてみましょう。

int var;

DPC++ は、上記の変数宣言を次のように変換します。

VarDecl  var 'int'

OpenCL* コンパイラーは、同じ変数宣言を次のように変換します。

VarDecl  var '__private int'

変数の型を変更すると、C++ では大規模で破壊的な影響があります。例えば、OpenCL* モードの C++ ではコンパイルできません。

template&lt;typename T1, typename T2&gt;
struct is_same {
    static constexpr int value = 0;
};

template&lt;typename T&gt;
struct is_same&lt;T, T&gt; {
    static constexpr int value = 1;
};

void foo(int p) {
    static_assert(is_same&lt;decltype(p), int&gt;::value, &quot;int is not an int?&quot;); // 失敗: p は '__private int' != 'int'
    static_assert(is_same&lt;decltype(&amp;p), int*&gt;::value, &quot;int* is not an int*?&quot;);  // 失敗: p は '__private int*' != '__generic int*'
}

既存の Clang の機能を利用するため、SYCL* モードでは以下の OpenCL* アドレス空間属性を再利用しています。

アドレス空間属性 SYCL* アドレス空間の列挙型
__attribute__((opencl_global)) global_space、constant_space
__attribute__((opencl_global_host)) ext_intel_global_host_space
__attribute__((opencl_global_device)) ext_intel_global_device_space
__attribute__((opencl_local)) local_space
__attribute__((opencl_private)) private_space
__attribute__((opencl_constant)) N/A

注: SYCL* デバイス・コンパイラーは __attribute__((opencl_constant)) をサポートしていますが、この属性の使用は SYCL* 実装内に制限されます。OpenCL* 定数ポインターは、ほかのアドレス空間 (デフォルトを含む) を持つポインターにキャストできません。

コンパイラー/ランタイム・インターフェイス

DPC++ ランタイム・アーキテクチャー

TBD

DPC++ 言語の SYCL* 拡張

言語拡張に一覧はこちら (英語) を参照してください。

タイトルとURLをコピーしました