インテル® oneAPI
CUDA
GPU
NVIDIA

oneAPI for NVIDIA* GPU 2024.0.0 ガイド

同カテゴリーの次の記事

ベータ版 oneAPI for AMD* GPU 2024.0.0 ガイド

この記事は、Codeplay 社の許可を得て iSUS (IA Software User Society) が作成した 2023年11月24日時点の『oneAPI for NVIDIA® GPUs 2024.0.0』の日本語参考訳です。原文は更新される可能性があります。原文と翻訳文の内容が異なる場合は原文を優先してください。

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

以前のバージョンのガイド: 2023.0.0 | 2023.1.0 | 2023.2.1
『ベータ版 oneAPI for AMD* GPU ガイド』はこちら

oneAPI for NVIDIA* GPU は、開発者が DPC++/SYCL* を利用して oneAPI アプリケーションを作成し、それらを NVIDIA* GPU 上で実行できるようにするインテル® oneAPI ツールキット向けのプラグインです。

このプラグインは、CUDA* バックエンドを DPC++ 環境に追加します。このドキュメントでは、「oneAPI for NVIDIA* GPU」と「DPC++ CUDA* プラグイン」は同じ意味で使われています。

oneAPI の詳細については、インテル® oneAPI の概要 (英語) を参照してください。

oneAPI for NVIDIA* GPU の使用を開始するには、「導入ガイド」を参照ください。

導入ガイド

パフォーマンス・ガイド

サポート

導入ガイド

oneAPI for NVIDIA* GPU のインストール

このガイドでは、DPC++ と DPC++ CUDA* プラグインを使用して、NVIDIA* GPU で SYCL* アプリケーションを実行する方法を説明します。

DPC++ に関連する一般的な情報は、「DPC++ のリソース」の節を参照してください。

サポートされるプラットフォーム

このリリースは、次のプラットフォームで検証されています。

GPU ハードウェア アーキテクチャー オペレーティング・システム CUDA* GPU ドライバー
NVIDIA* A100 PCIe* 40GB Ampere – sm_80 Ubuntu* 22.04.2 LTS 12.2 535.54.03
  • このリリースは各種 NVIDIA* GPU と CUDA* バージョンで動作するはずですが、Codeplay は評価されていないプラットフォームでの正常な動作を保証するものではありません。
  • このパッケージは Ubuntu* 22.04 でのみテストされていますが、一般的な Linux* システムにインストールできます。
  • このリリースの oneAPI for NVIDIA* GPU プラグインは、Windows* 上の oneAPI では使用できませんが、Windows* 向けのパッケージは今後リリースする予定です。
  • プラグインは、システムにインストールされている CUDA* のバージョンに依存します。CUDA* が macOS* をサポートしなくなったため、oneAPI for NVIDIA* GPU パッケージは macOS* では利用できません。

要件

  1. C++ 開発ツールインストールします。

    oneAPI アプリケーションをビルドして実行するには、C++ 開発ツールの cmake、gcc、g++、make および pkg-config をインストールする必要があります。

    次のコンソールコマンドは、一般的な Linux* ディストリビューションに上記のツールをインストールします。

    Ubuntu*

    $ sudo apt update
$ sudo apt -y install cmake pkg-config build-essential

    Red Hat* と Fedora*

    $ sudo yum update
$ sudo yum -y install cmake pkgconfig
$ sudo yum groupinstall "Development Tools"

    SUSE*

    $ sudo zypper update
$ sudo zypper --non-interactive install cmake pkg-config
$ sudo zypper --non-interactive install pattern devel_C_C++

    次のコマンドで、ツールがインストールされていることを確認します。

    $ which cmake pkg-config make gcc g++

    次のような出力が得られるはずです。

    /usr/bin/cmake
/usr/bin/pkg-config
/usr/bin/make
/usr/bin/gcc
/usr/bin/g++

  2. DPC++/C++ コンパイラーを含むインテル® oneAPI ツールキット 2024.0.0 をインストールします。

    • インテル® oneAPI ベース・ツールキットは、多くの利用環境に適用できます。

    • oneAPI for NVIDIA* GPU をインストールするには、インテル® oneAPI ツールキットのバージョン 2024.0.0 が必要です。これよりも古いバージョンにはインストールできません。

  3. 「Linux* 向けの NVIDIA* CUDA* インストール・ガイド」 (英語) の手順に従って、NVIDIA* GPU ドライバーと CUDA* ソフトウェア・スタックをインストールします。

インストール

  1. oneAPI for NVIDIA* GPU のインストーラー (英語) をダウンロードします。

  2. ダウンロードした自己展開型インストーラーを実行します。

    $ sh oneapi-for-nvidia-gpus-2024.0.0-cuda-12.0-linux.sh

    • インストーラーは、デフォルトの場所にあるインテル® oneAPI ツールキット 2024.0.0 のインストールを検索します。インテル® oneAPI ツールキットが独自の場所にインストールされている場合、--install-dir /path/to/intel/oneapi でパスを指定します。

    • インテル® oneAPI ツールキットが home ディレクトリー外にある場合、sudo を使用してコマンドを実行する必要があります。

環境を設定

  1. 実行中のセッションで oneAPI 環境を設定するには、インテルが提供する setvars.sh スクリプトを source します。

    システム全体へのインストールの場合:

    $ . /opt/intel/oneapi/setvars.sh --include-intel-llvm

    プライベート・インストールの場合 (デフォルトの場所):

    $ . ~/intel/oneapi/setvars.sh --include-intel-llvm

    • clang++ などの LLVM ツールにパスを追加するには、--include-intel-llvm オプションを使用します。

    • ターミナルを開くたびにこのスクリプトを実行する必要があります。セッションごとに設定を自動化する方法については、「CLI 開発向けの環境変数を設定する」 (英語) など、関連するインテル® oneAPI ツールキットのドキュメントを参照してください。

  2. CUDA* ライブラリーとツールが環境内にあることを確認します。

    1. nvidia-smi を実行します。実行時の表示に明らかなエラーが認められなければ、環境は正しく設定されています。

    2. 問題があれば、環境変数を手動で設定します。

      $ export PATH=/PATH_TO_CUDA_ROOT/bin:$PATH
$ export LD_LIBRARY_PATH=/PATH_TO_CUDA_ROOT/lib:$LD_LIBRARY_PATH

インストールの確認

DPC++ CUDA* プラグインのインストールを確認するには、DPC++ の sycl-ls ツールを使用して、SYCL* で利用可能な NVIDIA* GPU があることを確認します。NVIDIA* GPU が利用できる場合、 sycl-ls の出力に次のような情報が表示されます。

$ [ext_oneapi_cuda:gpu:0] NVIDIA CUDA BACKEND, NVIDIA A100-PCIE-40GB 8.8 [CUDA 12.2]

  • 上記のように利用可能な NVIDIA* GPU が表示されていれば、DPC++ CUDA* プラグインが適切にインストールされ、設定されていることが確認できます。

  • インストールや設定に問題がある場合、「トラブルシューティング」の「sycl-ls の出力でデバイスが見つからない場合」を確認してください。

  • 利用可能なハードウェアとインストールされている DPC++ プラグインに応じて、OpenCL* デバイス、インテル® GPU、または AMD* GPU など、ほかのデバイスもリストされることがあります。

サンプル・アプリケーションを実行

  1. 次の C++/SYCL* コードで構成される simple-sycl-app.cpp ファイルを作成します。

    #include <sycl/sycl.hpp>

int main() {
  // カーネルコード内で使用する 4 つの int バッファーを作成
  sycl::buffer<int, 1> Buffer{4};

  // SYCL* キューを作成
  sycl::queue Queue{};

  // カーネルのインデックス空間サイズ
  sycl::range<1> NumOfWorkItems{Buffer.size()};

  // キューへコマンドグループ (ワーク) を送信
  Queue.submit([&](sycl::handler &cgh) {

    // デバイス上のバッファーへの書き込み専用アクセサーを作成
    auto Accessor = Buffer.get_access<sycl::access::mode::write>(cgh);

    // カーネルを実行
    cgh.parallel_for<class FillBuffer>(
        NumOfWorkItems, [=](sycl::id<1> WIid) {
          // インデックスでバッファーを埋めます
          Accessor[WIid] = static_cast<int>(WIid.get(0));
        });
  });

  // ホスト上のバッファーへの読み取り専用アクセサーを作成。
  // キューのワークが完了するのを待機する暗黙のバリア
  auto HostAccessor = Buffer.get_host_access();

  // 結果をチェック
  bool MismatchFound{false};
  for (size_t I{0}; I < Buffer.size(); ++I) {
    if (HostAccessor[I] != I) {
      std::cout << "The result is incorrect for element: " << I
                << " , expected: " << I << " , got: " << HostAccessor[I]
                << std::endl;
      MismatchFound = true;
    }
  }

  if (!MismatchFound) {
    std::cout << "The results are correct!" << std::endl;
  }

  return MismatchFound;
}

  2. アプリケーションをコンパイルします。

    $ icpx -fsycl -fsycl-targets=nvptx64-nvidia-cuda simple-sycl-app.cpp -o simple-sycl-app

    インストールされている CUDA* のバージョンによっては、次のような警告が表示されることがありますが、これは無視してもかまいません。

    $ icpx: warning: CUDA version is newer than the latest supported version 12.1 [-Wunknown-cuda-version]

  3. アプリケーションを実行します。

    $ ONEAPI_DEVICE_SELECTOR="ext_oneapi_cuda:*" SYCL_PI_TRACE=1 ./simple-sycl-app

    次のような出力が得られます。

    SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_cuda.so [ PluginVersion: 14.37.1 ]
SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_unified_runtime.so [ PluginVersion: 14.37.1 ]
SYCL_PI_TRACE[all]: Requested device_type: info::device_type::automatic
SYCL_PI_TRACE[all]: Selected device: -> final score = 1500
SYCL_PI_TRACE[all]:   platform: NVIDIA CUDA BACKEND
SYCL_PI_TRACE[all]:   device: NVIDIA A100-PCIE-40GB
The results are correct!

    これで、oneAPI for NVIDIA* GPU の環境設定が確認でき、oneAPI アプリケーションの開発を開始できます。

    以降では、NVIDIA* GPU で oneAPI アプリケーションをコンパイルして実行するための一般的な情報を説明します。

DPC++ を使用して NVIDIA* GPU をターゲットにする

NVIDIA* GPU 向けのコンパイル

NVIDIA* GPU 対応の SYCL* アプリケーションをコンパイルするには、DPC++ に含まれる icpx コンパイラーを使用します。

例:

$ icpx -fsycl -fsycl-targets=nvptx64-nvidia-cuda sycl-app.cpp -o sycl-app

次のフラグが必要です。

  • -fsycl: C++ ソースファイルを SYCL* モードでコンパイルするようにコンパイラーに指示します。このフラグは暗黙的に C++ 17 を有効にし、SYCL* ランタイム・ライブラリーを自動でリンクします。

  • -fsycl-targets=nvptx64-nvidia-cuda: NVIDIA* GPU をターゲットとして、SYCL* カーネルをビルドすることをコンパイラーに指示します。

また、次のフラグを使用して、特定の NVIDIA* アーキテクチャー向けの SYCL* カーネルをビルドすることができます。

  • -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_80

デフォルトではカーネルは sm_50 用にビルドされ、多様なアーキテクチャーで動作しますが、新しい CUDA* 機能の利用は制限されることに注意してください。

利用できる SYCL* コンパイルフラグの詳細は、『DPC++ コンパイラー・ユーザーズ・マニュアル』 (英語) を参照してください。すべての DPC++ コンパイラー・オプションの詳細は、『インテル® oneAPI DPC++/C++ コンパイラー・デベロッパー・ガイドおよびリファレンス』の「コンパイラー・オプション」 (英語) を参照してください。

icpx コンパイラーを使用する

icpx コンパイラーは、デフォルトで -O2 と -ffast-math オプションを有効にするため、通常の clang++ ドライバーよりも積極的な最適化を行います。多くの場合、これによりパフォーマンスは向上しますが、一部のアプリケーションでは問題が生じる可能性があります。その場合、-fno-fast-math を使用して -ffast-math を無効にして、-O2 以外の -O オプションを指定することで最適化レベルを変更できます。$releasedir/compiler/latest/linux/bin-llvm/clang++ にある clang++ ドライバーを直接起動することで、通常の clang++ の最適化レベルを適用できます。

複数ターゲット向けのコンパイル

NVIDIA* GPU をターゲットにするだけではなく、一度のコンパイルで複数のハードウェア・ターゲットで実行できる SYCL* アプリケーションを生成できます。次の例は、NVIDIA* GPU、AMD* GPU、および SPIR* をサポートする任意のデバイス (インテル® GPU など) で実行できるコードを含む単一のバイナリーを生成する方法を示しています。

$ icpx -fsycl -fsycl-targets=amdgcn-amd-amdhsa,nvptx64-nvidia-cuda,spir64 \
        -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1030 \
        -Xsycl-target-backend=nvptx64-nvidia-cuda --offload-arch=sm_80 \
        -o sycl-app sycl-app.cpp

NVIDIA* GPU でアプリケーションを実行

NVIDIA* ターゲット向けに SYCL* アプリケーションをコンパイルしたら、ランタイムが SYCL* デバイスとして NVIDIA* GPU を選択しているか確認する必要があります。

通常、デフォルトのデバイスセレクターを使用するだけで、利用可能な NVIDIA* GPU の 1 つが選択されます。しかし、場合によっては、SYCL* アプリケーションを変更して、GPU セレクターやカスタムセレクターなど、より正確な SYCL* デバイスセレクターを設定することもあります。

環境変数 ONEAPI_DEVICE_SELECTOR を設定して、利用可能なデバイスセットを限定することで SYCL* デバイスセレクターを支援できます。例えば、DPC++ CUDA* プラグインでサポートされるデバイスのみを許可するには、次のように設定します。

$ export ONEAPI_DEVICE_SELECTOR="ext_oneapi_cuda:*"

この環境変数の詳細については、インテル® oneAPI DPC++ コンパイラーのドキュメントで「環境変数」 (英語) 参照してください。

DPC++ のリソース

SYCL* のリソース

 

SYCL* アプリケーションのデバッグ

この節では、さまざまなデバイスで SYCL* アプリケーションをデバッグするための情報、ヒント、およびポインターについて説明します。

SYCL* アプリケーションのホストコードは、単純に C++ アプリケーションとしてデバッグできますが、カーネルデバッグのサポートやツールは、ターゲットデバイスによって異なる可能性があります。

注意
SYCL* アプリケーションに汎用性がある場合、実際のターゲットデバイスではなく、インテルの OpenCL* CPU デバイスなど、豊富なデバッグサポートとツールを備えたデバイスでデバッグしたほうが有用なことがあります。

インテルの OpenCL* CPU デバイスでのデバッグ

インテルの OpenCL* CPU デバイスを使用した DPC++ アプリケーションのデバッグについては、『インテル® oneAPI プログラミング・ガイド』の「DPC++ と OpenMP* オフロードプロセスのデバッグ」の節を参照してください。

CUDA* デバッガーのサポート

CUDA* ツールキットには、CUDA* アプリケーションの NVIDIA* GPU カーネルのデバッグをサポートする cuda-gdb デバッガーが付属しています。cuda-gdb は、DPC++ ext_oneapi_cuda バックエンド用にコンパイルされたカーネルのデバッグにも使用できます。現在、NVIDIA* nvcc コンパイラーでコンパイルされたカーネルと比較して、DPC++ でコンパイルされたカーネルのデバッグに cuda-gdb が利用される場合、予測される動作に違いは報告されていません。cuda-gdb の詳しい使用法については、cuda-gdb のドキュメント (英語) を参照してください。

CUDA* バックエンドの MPI ガイド

このセクションでは、DPC++ CUDA* バックエンドで CUDA* 対応の MPI を使用する方法を説明します。

必要要件

ここでは、CUDA* バックエンドをサポートするインテル® oneAPI DPC++ コンパイラーが正常にインストール済みであることを前提としています。DPC++ のインストールと設定方法については、導入ガイドを参照してください。また、DPC++ コンパイラーを起動するコンパイラー・ラッパー (mpicxx など) を使用する CUDA* 対応の MPI 実装も必要になります。

CUDA バックエンドで MPI を使用する

アプリケーションのコンパイルと実行

send_recv_buff.cpp および send_recv_usm.cpp サンプルコードは、バッファーまたは USM を使用して CUDA* 対応 MPI を DPC++ で使用する例を示す入門サンプルです。CUDA* 対応 MPI でバッファーを使用するには、host_task 内で MPI 呼び出しを行う必要があります。詳細については、send_recv_buff.cpp を参照してください。 MPI で SYCL* USM を使用するには、常にメインスレッドから MPI 関数を直接呼び出す必要があります。host_tack 内から SYCL* USM を取得する MPI 関数の呼び出しの動作は、現在未定義です。さらに、send_recv_buff.cpp のサンプルでは、MPI を SYCL* 2020 のリダクションおよび parallel_for と併用して、最適化された簡略な複数ランクでのリダクションを実現する方法を示しています。

サンプルをコンパイルするには、コンパイラー・ラッパー (mpicxx など) が DPC++ コンパイラーを起動できるようにする必要があります。最初に、パスにラッパーが設定されていることを確認します。

$ export PATH=/path/to/your-mpi-install/bin:$PATH

次に、サンプルをコンパイルします。

$ mpicxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_xx send_recv_usm.cpp -o ./res

ここで、「sm_xx」は、デバイスのアーキテクチャーを指定します。

サンプルを実行するには 2 つのランクが必要です。次のコマンドで MPI を使用してアプリケーションを実行します。

$ mpirun -n 2 ./res

「-n 2」は、2 つのランクを使用してアプリケーションを実行することを指定します。サンプルに変更を加えるか、各ランクで特別な環境変数を設定しない限り、これは 1 つの GPU を使用して 2 つのランクが実行されることに注意してください。詳しくは、次のセクションで説明します。

特定の GPU に MPI ランクを割り当て

単一ノード内の CUDA* 対応 MPI では、ユーザーは各ランクが特定の GPU を使用するかどうかを制御する機能が必要になります。これを実現する 1 つの方法は、プログラム内で特定のランクを特定の GPU にマップすることです。現在、ext_oneapi_cuda バックエンドでは、各 CUDA* デバイスに独自のプラットフォームを持ちます。sycl::platform::get_platforms() から返されるプラットフォームのベクトルは、CUDA* デバイス ID で順序付けされ、利用可能な最小 ID が最初に配置されます。これは、デバイス 0 と 1 が利用できる場合、次の方法でキューを初期化することで、2 ランクの MPI プログラムのランク 0 と 1 にこれらのデバイスを割り当てることを意味します。

std::vector<sycl::device> Devs;

for (const auto &plt : sycl::platform::get_platforms()) {
  if (plt.get_backend() == sycl::backend::ext_oneapi_cuda)
    Devs.push_back(plt.get_devices()[0]);
}

sycl::queue q{Devs[rank]};

注: CUDA* デバイスは単一のプラットフォームに含まれるため、上記のコードを調整する必要があります。NVIDIA* GPU のみが利用可能である場合、よりシンプルな方法でリストを作成できます。

std::vector<sycl::device> Devs =
      sycl::device::get_devices(sycl::info::device_type::gpu);

MPI ランクを一意の GPU にマッピングする別の方法は、MPI プロセスごとに CUDA_VISIBLE_DEVICES 環境変数を単一の値にすることです。一般的には、CUDA_VISIBLE_DEVICES をローカル MPI ランク ID と同じ値に設定します。ローカルランク ID を取得する方法は、MPI 実装のドキュメントを参照してください。

SLURM システムを利用する場合、GPU アフィニティー・オプション -gpu-bind を使用して、CUDA_VISIBLE_DEVICES と同じ効果を実現できます。詳細については、SLURM のドキュメントを参照してください。

制限事項

  • DPC++ を使用する CUDA* 対応の MPI では、現在ノード間の MPI の SYCL* 共有 USM をサポートしていません。

 

パフォーマンス・ガイド

はじめに

このガイドは、SYCL* プログラミング・モデルと一般的な GPU におけるパフォーマンスの紹介から始まります。次に、GPU でのパフォーマンス解析の基本と、そこで使用される一般的なツールを紹介します。最後に、ベンダー固有の GPU と利用可能なツールについて紹介します。また、無料の書籍『Data Parallel C++』 (英語) を一読されることを推奨します。第 15 章では、SYCL* と DPC++ に関連した GPU 上でのパフォーマンスについて説明されています。

NVIDIA* GPU と AMD* GPU の両方に適用される一般的な SYCL* 最適化については、「一般的な最適化」を参照してください。

NVIDIA* GPU をターゲットにする固有の最適化については、「NVIDIA* GPU 上のパフォーマンス」を参照してください。

インテル® GPU 固有のパフォーマンス最適化については、対応するインテル® GPU 固有のパフォーマンス・ガイドを参照してください。

プログラミング・モデル

グラフィックス処理ユニットは、超並列アーキテクチャーにより、CPU よりも 1 秒あたり多くの浮動小数点演算を実行でき、メモリー帯域幅も高くなっています。これらの機能は、コードの開発時点で GPU アーキテクチャーを使用することを選択した場合にのみ活用できます。

ここでは、GPU における大規模並列処理を表現するプログラミング・モデルが基本となります。SYCL* は OpenCL* や CUDA* と同様のプログラミング・モデルを採用しており、カーネル (GPU によって実行される関数) は work-item によって実行される操作で表現されます。

SYCL* 仕様 (Rev 8) の 3.7.2 節 (英語) では次のように定義されています。

カーネルが実行のため送信されると、インデックス空間が定義されます。カーネルボディーのインスタンスは、インデックス空間の各ポイントで実行されます。カーネル・インスタンスは work-item (ワーク項目) と呼ばれ、グローバル id を提供するインデックス空間内のポイントで識別されます。それぞれの work-item は同じコードを実行しますが、コードと操作されるデータの実行パスは、work-item のグローバル id を使用して計算を特殊化することで異なります。

SYCL* では、2 つの異なるカーネル実行モデルを利用できます。

SYCL* 仕様 (Rev 8) の 3.7.2.1 節 (英語) では次のように記述されています。

range<N> (N は 1、2 または 3) で定義される N 次元のインデックス空間でカーネルを呼び出す単純な実行モデルをサポートします。この場合、カーネルの work-item は独立して実行されます。各 work-item は、タイプ item<N> の値によって識別されます。タイプ item<N> は、タイプ id<N> の work-item 識別子と、カーネルを実行する work-item の数を示す range<N> をカプセル化します。

SYCL* 仕様 (Rev 8) の 3.7.2.2 節 (英語) では次のように記述されています。

work-item を wark-group に編成できる ND-range の実行モデルは インデックス空間よりも粗い粒度の分解を提供します。それぞれの work-group には、work-item で使用できるインデックス空間と同じ次元の work-group id が割り当てられます。work-item には、それぞれ work-group 内で一意のローカル id が割り当てられるため、単一 work-item は、グローバル id、またはローカル id と work-group id の組み合わせで識別できます。特定の work-group 内の work-item は、単一の計算ユニットの処理ユニットで同時に実行されます。SYCL* で使用される work-group は、ND-range と呼ばれます。ND-range は、N 次元のインデックス空間であり、N は 1、2 または 3 です。SYCL* では、ND-range は nd_range<N> クラスを介して表現されます。nd_range<N> は、グローバルレンジとローカルレンジで構成され、それぞれ range<N> タイプの値で表現されます。さらに、タイプ id<N> 値で表現されるグローバルオフセットが存在することもありますが、これは SYCL* 2020 では非推奨です。タイプ nd_range<N> と id<N> は、それぞれ N 要素の整数配列です。nd_range<N> で定義される反復回数は、ND-range のグローバルオフセットで開始される N 次元のインデックス空間であり、サイズはグローバルレンジで、ローカル・レンジ・サイズの work-group に分割されます。ND-range の各 work-item は、タイプ nd_range<N> の値によって識別されます。タイプ nd_range<N> は、グローバル id、ローカル id、および work-group id をすべて id<N> (id<N> タイプの反復空間オフセットですが、SYCL* 2020 では非推奨) としてカプセル化し、グローバルとローカルレンジを同期して work-group を有効にします。work-group には、work-item のグローバル id と同様の方法で id が割り当てられます。work-item には work-group とゼロからその次元の work-group サイズから 1 を引いた範囲のコンポーネントを保持するローカル id が割り当てられます。つまり、work-group id と work-group 内のローカル id の組み合わせで work-item が一意に定義されます。

work-item は、次の OpenCL* メモリーモデルに従って 3 つの異なるメモリー領域にアクセスできます。

  • グローバルメモリー: すべての work-group のすべての work-item 間で共有されます。

  • ローカルメモリー: 同一 work-group のすべての work-item 間で共有されます。

  • プライベート・メモリー: 各 work-item でプライベートです。

アーキテクチャー

SYCL* 仕様では、独立して動作する 1 つ以上の計算ユニット (CU) で構成されるデバイスを考慮することで、OpenCL* 1.2 の仕様に従います。NVIDIA では CU を ストリーミング・マルチプロセッサー (streaming multiprocessor) と呼び、AMD では単純に計算ユニット (compute unit) と呼んでいます。それぞれの CU は、1 つ以上の処理エレメント (PE) とローカルメモリーで構成されます。work-group は単一の CU で実行されますが、work-item は 1 つ以上の PE で実行されることがあります。一般に、CU は SIMD 形式で work-item の小さなセット (sub-group として定義) を実行します。sub-group は NVIDIA では ワープ (warp) 、AMD ではウェーブフロント (wavefront) と呼ばれます。sub-group サイズは NVIDIA 向けには 32 で、AMD 向けには通常 64 (一部のアーキテクチャー向けには 32) です。

計算

カーネルを構成する work-group は、CU 全体にスケジュールされます。この時点で、それぞれの CU は処理エレメントで 1 つ以上の sub-group を実行します。計算ユニットには、算術演算を実行する整数論理ユニットや浮動小数点ユニット、メモリー操作を行うロード/ストアユニット、超越関数 (正弦、余弦、逆数、平方根など) を実行する特別なユニット、AI で役立つ行列操作など、さまざまな種類の処理エレメントが含まれます。処理エレメントが操作を完了するのに要する時間 (クロックサイクルで測定) は、レイテンシーと呼ばれます。レイテンシーは操作の種類によって異なります。例えば、グローバル・メモリー・トランザクションのレイテンシーは、レジスター呼び出しに比べ桁違いに大きく、これは各種算術演算でも同じことが当てはまります。

スループットは、実行された操作の数と、それらの完了に要する時間の比率です。この比率は、命令のレイテンシーを減らすか、同時に実行する命令数を増やすことで高めることができます。これまで、CPU はクロック周波数を上げて命令レイテンシーを最小化することでスループットを向上させてきました。一方、GPU はレイテンシーを隠匿することでスループットを向上させます。これにより、CU は sub-group 間で「コンテキスト」 (レジスター、命令カウンターなど) をわずかな労力で変更できます。そのため、操作に多くのクロックサイクルを要する場合、CU は「コンテキスト」を変更し、別の sub-group の操作を実行することでそれらを隠匿できます。アーキテクチャーによって、同時に実行できる sub-group の最大数は異なります。実際に実行中の sub-group と実行中の sub-group の最大数の比率は「占有率」として定義されます。次の節で詳しく説明します。

GPU における work-item の同時実行は、複数レベルで実現されます。

  1. 同一 sub-group 内の異なる work-item は SIMD 形式で同期実行されます。つまり、同じ操作が異なるデータを実行します。

  2. 前述したように、CU はレイテンシーを隠匿するため、同一または異なる work-group から複数の sub-group を同時に実行します。

  3. GPU を構成する CU は、異なる work-group に属する、異なる sub-group を同時に実行します。

これらの並列実行の機能は、起動されたカーネルが GPU 全体をビジー状態にする十分な大きさの work-item を持っている場合にフル活用されます。

メモリー

次の図は、ディスクリート GPU を搭載したシステムにおける一般的な接続方法を示しています。[1] ホストとデバイスを接続し、[2] CU をグローバルメモリーに接続します。例えば、NVIDIA* GA100 GPU の目安となる帯域幅は次のようになります。[1] PCIe* x16 4.0 では 31GB/秒、および [2] HBM2 では 1555GB/秒。


図 1

CPU と GPU 間の接続 [1] が大きなボトルネックになる可能性があります。そのため、ホストとデバイス間のデータ転送を慎重に検討し、GPU 上のデータの局所性を可能な限り維持することが重要です。ただし、カーネルの実行とオーバーラップすることで、PCIe* メモリーのトランザクションで生じるレイテンシーを隠匿することができます。

GPU の主要な特徴として、CU とグローバルメモリー間の高い帯域幅があります [3]。これは、それらを接続するメモリー・コントローラーの数と幅によるものです。例えば、NVIDIA* GA100 GPU には、12 個の 512 ビットの HBM メモリー・コントローラーがあります。これにより、クロックサイクルごとに大量のデータを転送できます。NVIDIA* GA100 GPU では、クロックごとに 6144 ビットです。ただし、この高帯域幅のメモリーを十分に活用するには、メモリーアクセスを結合する必要があります。つまり、work-item はキャッシュに最適な方法でメモリーアクセスする必要があります。


図 2

work-item とグローバルメモリー間にはいくつかのメモリー階層があります。以下に、それらをアクセス・レイテンシーが低い順に示します。

  • レジスターは、ワークメモリーとして使用される work-item からは透過なデータを維持します。

  • コンスタント・メモリーは、CU が使用する読み取り専用メモリーです。

  • ローカルメモリーは CU ごとにあり、同一 work-group 内の work-item 間で共有されます。ローカルメモリーは、グローバルメモリーよりも高速であり、再利用されるグローバルメモリーのデータをキャッシュするために使用されます。

  • グローバルメモリー (DDR または HBM) と CU を接続するメモリーシステムを構成する L1 および L2 キャッシュ。

最適化の目的

優先度

GPU コードのパフォーマンスに影響する主な要因を重要度の高い順に示します。

  • 合成されていないグローバル・メモリー・アクセス。キャッシュが完全に活用されると、メモリーアクセスは結合され、高い帯域幅を維持できます。結合の方法はアーキテクチャーによって異なりますが、一般に、同じ sub-group 内の work-item が連続したメモリー位置をアクセスすることで実現されます。

  • ローカルメモリーのバンク競合。ローカルメモリーは複数のバンクに分割されており、異なる work-item から同時にアクセスできます。異なる work-item が同じメモリーバンクにアクセスすると、バンク競合が発生してトランザクションはシリアル化されます。

  • if 文などの条件式やループの反復回数は work-item によって異なるため、同じ sub-group に属する work-item が異なる命令を実行することで発散が発生します。近年のアーキテクチャーではこの事象が緩和され、パフォーマンスのペナルティーが軽減されています。

計算の種類が異なれば最適化の優先順位も変わってきます。例えば、メモリー・トランザクションに対し算術演算が少ないメモリー依存のタスクを考えてみます。この場合、GPU を十分に活用するには、メモリーアクセスを結合することが重要です。一方、メモリー・トランザクションに対し算術演算が多い計算依存タスクがあります。この場合、スレッドの発散を回避することが有用な場合があります。算術演算数とリード/ライトデータのバイト数の比率は、演算強度として定義されます。

I = (浮動小数点操作数) / (リード/ライトデータのバイト数) [FLOP/バイト]

ルーフライン・モデルを利用して、カーネルの演算強度をハードウェア特性に関連付けることで、カーネルがメモリー依存であるか計算依存であるか確認できます。ルーフライン・モデルは 2 次元座標として表示され、x 軸には演算強度が、y 軸には浮動小数点演算のスループット (FLOPS: 1 秒あたりの浮動小数点演算) が示されます。


図 3

実際のルーフラインを構成する最初のセグメントは、y = x * B を示します。ここで、B はグローバル・メモリー・システムの帯域幅です。次に水平線 (y  = Pmax) は、FMA など特定の操作の最大浮動小数点スループット (Pmax) に依存します。セグメントが遭遇するポイントはリッジポイントと呼ばれます。

カーネルのパフォーマンスは、ルーフラインにプロットされたポイント (点) で示されます。x 軸はカーネルの演算強度を示し、y 軸は計測されたカーネルの FLOPS を示します。このポイントがリッジポイントの左にある場合、そのカーネルはメモリー依存であり、右にある場合は計算依存です。

占有率

カーネルのパフォーマンスを評価するには、その占有率を考慮します。占有率は、次のように定義される計算ユニットの式で求められます。

占有率 = アクティブな sub-group 数 / アクティブな sub-group の最大数

アクティブな sub-group は、CU で実際に実行される sub-group です。アクティブな sub-group の最大数は、計算ユニットのアーキテクチャーによって異なります。例えば、NVIDIA* GA100 CU アーキテクチャーでは 64 です。

占有率を高めるにはアクティブな sub-group 数を最大化する必要があります。ただし、計算ユニットのアーキテクチャーによって制約は異なります。

  • work-group あたりの work-item の最大数

  • CU で実行される work-group の最大数: work-group サイズが小さすぎると、CU はアクティブな sub-group の最大数を実行することができません。

  • レジスター数の制限: カーネルコードが複雑になるとレジスターの使用量が増加します。コードを簡素にすることでレジスターの使用量を軽減できます。これは、コードを複数のカーネルに分割することで実現することもできます。

  • ローカルメモリー量の制限: work-group がローカルメモリーを消費しすぎると、同時に実行できる work-group 数が減少します。

work-group が使用するレジスターやローカルメモリーが多いと、占有率が制限される可能性があります。ユーザーは、work-group のサイズを変更することで占有率を改善できます。このサイズは、sub-group サイズの倍数で、アクティブな sub-group の最大数の除数である必要があります。

例えば、NVIDIA* GA100 GPU では、各 work-item は最大 32 個のレジスターを使用して完全な占有を実現できます。

r_max = (CU ごとのレジスター数) / (アクティブな sub-group の最大数) * (sub-group サイズ) = 32

work-item が 32 未満のレジスターを使用する場合、CU で同時に実行できる work-group の最大数 wg_max とすると、ローカルメモリーにも同じことが当てはまります。

wg_max = (アクティブな sub-group の最大数) * (sub-group サイズ) / (実際の work-group サイズ)

各 work-group が最大 48Kb / wg_max のローカルメモリーを割り当てる場合、完全な占有率が得られます。

実効占有率は重要ですが、パフォーマンスにおける最重要のメトリックではありません。命令レベルの並列処理が十分にあり、同じ sub-group に属する独立した命令の同時実行が可能であれば、低い占有率でもレイテンシーを十分に隠匿できます。これについては、こちら (英語) をご覧ください。

さらに、GPU のすべての CU を利用するためカーネルで起動される work-item の最小数は、少なくとも次の wi_min でなければなりません。

wi_min = (アクティブな sub-group の最大数) * (sub-group サイズ) * (CU 数) = 262144

これらのパラメーターはすべて、特定ベンダーの各アーキテクチャーのドキュメントに記載されていますが、以下の表にいくつかの一般的な GPU アーキテクチャーの数値を示します。

一般的な GPU アーキテクチャーのリファレンス占有率
アーキテクチャー アクティブな
sub-group の最大数		 work-item の
最大数		 work-group の
最大数		 レジスター数 ローカルメモリー
(バイト)
NVIDIA* S.M. 7.0 64 2048 32 65536 65536
NVIDIA* S.M. 7.5 32 1024 16 65536 65536
NVIDIA* S.M. 8.0 64 2048 32 65536 65536
AMD* GFX9xx 40[1] 1024 16 29184 (?) 65536

 

[1] この図は、AMD アーキテクチャー全般に適用されます。work-group が 1 つの sub-group (例えば 64 work-item) のみである場合、CU あたりの work-group の最大数は 40 です。

NVIDIA* GPU の場合、NVIDIA* Nsight* Compute (英語) は占有計算機を提供しており、理論上の占有率がどのように計算されるか判断するのに役立ちます。非推奨ですが、NVIDIA の spreadsheet (英語) でも同様の機能を提供することを示しています。

パフォーマンス解析

パフォーマンス解析と最適化は繰り返し作業です。開発者は、ツールを使用してアプリケーションのパフォーマンスを測定しボトルネックを特定して、それらを改善しながら、この手順を繰り返します。それぞれの反復作業で、以前は見つからなかったボトルネックが明らかになることがあります。

ある時点で、アプリケーションの制限要因となる部分で、可能な限り高いパフォーマンスを特定することが重要です。これは、光速またはルーフラインと呼ばれることもあり、アプリケーションの理論上のピーク・パフォーマンスを予測したり、そのパフォーマンスにどれだけ近づいているかを判断するのに役立ちます。

以降の節では、解析ツールと制限要因について詳しく説明します。

解析の方法論

パフォーマンス解析に使用されるツールはプロファイラーとも呼ばれます。プロファイルという用語はいろいろな意味で使用されます。ここでは、パフォーマンス解析に使用されるツールの総称という意味で使用します。特定のパフォーマンス・ツールの説明では、より具体的な意味で使用されることがあります。

パフォーマンス解析は、大きく分けてトレースとサンプリングに分類されます。トレースは、アプリケーションの実行中に 1 つ以上のイベントが発生するたびに記録します。サンプリングは、実行中のアプリケーションの状態を定期的に調査して、その状態を記録します。頻繁に発生するイベントでは、トレースで大量のデータが蓄積される可能性があります。サンプリングでは、サンプリング間隔を調整することでデータ量を制御できます。間隔を長くするとデータ量は減りますが、短い間隔の動作を記録できないことがあります。どちらも改良すべき点はありますが、トレースまたはサンプリングのいずれかを実行中のデータ軽減と組み合わせることができます。

どの解析ツールでも、考慮すべき 2 つのことがあります。

  • オーバーヘッド: ツールが通常のプログラムの実行時間をどれくらい増加させるかを表わします。パフォーマンス・ツールは、オーバーヘッドを最小限に抑えることが求められます。ただし、オーバーヘッドの増加を十分に理解している場合は、データを解釈する際にこれを補正できます。例えば、あるツールはコードの GPU 実行領域では正確な結果を提供し、CPU 実行領域では実行時間が長くなります。

  • データ量: 生成されるファイルの大きさを示します。データ量が多いと、オーバーヘッドも増加します。また、大きな出力データセットは管理が困難で、特に出力データセットを表示するためリモートマシンに移動する場合、後処理ツールの応答性にも問題があります。

システムレベルの解析

システムレベルの解析では、同一ノードまたは異なるノード上のプロセス間の相互作用、および CPU と GPU 間の相互作用を調査します。

複雑なワークロードにおける CPU と GPU 間の相互作用を解析するのは困難なことがあります。ベンダーは、このような解析を支援するためトレースツールを提供することがあります。それらは、メモリー割り当て、メモリー転送、カーネルの起動、同期など、GPU 間の API 呼び出しのタイムスタンプと期間を記録します。これらのツールには、シリアル化や過度のアイドル時間などのボトルネックを視覚的に特定するタイムライン表示が含まれます。

状況に応じて、OS のカーネルトレース (Linux* ftrace など） を使用して、それをアプリケーションの実行に関連付けると便利です。これには、root 権限が必要になります。パフォーマンスの問題に関連するカーネルのアクティビティーが理解できない場合は、循環バッファーを利用するすべての OS アクティビティーを記録し、パフォーマンスの問題が検出されたときにアプリケーションの制御下でバッファーをダンプすると便利です (例えば、タイムステップが平均時間や予測時間よりも大幅に長くかかる場合)。循環バッファーによる手法は、トレース・データ・ストリーム全体を記録する際にコストが高い場合に有効です。

分散アプリケーションのスケーリング (通常メッセージ・パッシング・インターフェイス (英語) を使用) は特筆に値します。一般に使用されるスケーリングには 2 つの定義があります。強力なスケーリングは、問題のサイズを一定に保ち、MPI ランクの数が増加するのにしたがって経過時間を測定します。弱いスケーリングは、MPI ランクの数に比例して問題サイズを大きくします。

強力なスケーリングはより困難な問題です。多くの場合、すべての MPI ランクをビジーに保つのに十分なワークがありません。MPI プロファイル・ツールを利用して、異なる数のランクでスイープを実行し、MPI プロファイルを比較することが有用です。

特に大規模なスケーリングでは、そのほかの MPI の問題がしばしば発生します。リダクション操作は log(N) に反比例します。さらに、小さなリダクション操作 (スカラー値への MPI_Allreduce など) は、OS によるノイズの影響を受ける可能性があります。ネットワークが混雑する可能性があるため、大規模な共有クラスターではポイントツーポイント操作でも影響を受ける可能性があります。強力なスケーリングでは、メッセージサイズは通常、ランク数が多いほど小さくなるため、MPI レイテンシーがさらに重要になります。

開発者は、アプリケーションの動作が大規模なノードと小規模なノードで実行される際の違いを予測する必要があります。通常のように MPI プロファイル・ツールを使用すると、動作の違いを理解するのに役立ちます。オーバーヘッドの低いツールは、大規模なケースでは特に重要です。

カーネルレベルの解析

カーネルレベルの解析では、GPU カーネルの実行に費やされた時間と、個々の GPU カーネルのパフォーマンスに注目します。

前の節で説明したツールは、通常、起動パラメーター、起動回数、カーネルで消費された時間など、カーネル実行ごとのサマリーを示します。アプリケーションの合計時間は、CPU の経過時間と GPU の経過時間の合計として見積もられることが多く、GPU での経過時間はカーネルの実行時間の合計として概算されます。これにより、GPU カーネルの実行時間を改善することで、全体でどれだけ改善されるかが分かります。実行がオーバーラップしていたり、データの転送時間が長い場合は、常に正確であるとは限りませんが、経験則としては適切です。

カーネルのパフォーマンスを詳しく解析するには以下が必要です。

  • カーネルのソースコードを調査
  • カーネル向けにコンパイラーが生成するアセンブリー言語の調査
  • カーネル実行中のハードウェア・パフォーマンス・メトリックの収集

アセンブリー言語を生成する方法は、コンパイラーと GPU によって異なります。詳細については以降で説明します。

ここからは、GPU (多くの場合 CPU にも該当) で利用可能なメトリックと、それらを解釈してパフォーマンスを改善する作業で導入できる一般的な手法について説明します。異なる GPU 向けの詳細については、このドキュメントの後半で説明します。

重要な GPU メトリック

レートメトリック

アプリケーションが GPU を使用するのは、利用可能な計算リソースを増やすためです。通常、計算スループットは、単位時間あたりに処理される演算数で示されます。例えば、倍精度浮動小数点演算の数/秒、32 ビット整数演算の数/秒などです。特定の GPU では、これらのピーク値が公開されています。

多くの場合、アプリケーションのピーク・パフォーマンスは、非計算リソース (特にメインメモリーやスクラッチパッド・メモリーなどさまざまなメモリー領域) へのアクセスによって制限されます。ここにもピーク値があります。例えば、メインメモリーの帯域幅は、単位時間あたりのバイト数で表現されます。

従来のルーフラインのようなモデルでは、ほかのリソースによる制限 (一般的なものはメインメモリーの帯域幅) を考慮して、達成可能な計算パフォーマンスを定量化しようとします。アプリケーションが計算以外のメトリックでピーク・パフォーマンスに達している場合、ピーク計算パフォーマンスを達成することはできません。これにより、開発者は、アプリケーションで達成可能なピーク・パフォーマンスに関する情報を得ることができます。

利用率メトリック

特定のリソースや機能ユニットがどれだけビジーであるかを知るのは有用です。この利用率メトリックは、レートメトリックとは異なります。リソースの利用率が高くても、ピーク・パフォーマンスにほど遠い場合があります。1 つの例として、メモリー・アクセス・パターンが不均一なカーネルが挙げられます。この場合、メモリー帯域幅がピークから離れていても、メモリーユニットの利用率は非常に高くなることがあります。利用率メトリックは、ルーフライン・モデルでは明らかにならないボトルネックを理解するのに役立ちます。

通常、利用率メトリックはメモリーユニットと計算ユニットで利用できます。また、キャッシュやローカルメモリーなど、各種マイクロアーキテクチャー・ブロックでも利用できることがあります。

発散

前述のように、GPU は複数の work-item を SIMD (単一命令複数データ) 方式で同時に実行する複数の計算ユニット (CU) で構成されています。

開発者は、単一の work-item に対して実行する操作を記述します。コンパイラーは、このコードを複数の work-item を同時に処理する命令に変換します。各 GPU には、sub-group サイズと呼ばれる、同時に実行される work-item の最小数がネイティブに設定されています。

発散 (Divergence) は、異なる work-item が異なるパスをたどることで発生します。多くの work-item が特定の命令で実行される場合、コンパイラーは可能なすべてのパスの組み合わせを考慮して命令を生成する必要があります。特定の命令で非アクティブな work-item は無効になります。これにより SIMD レーンの一部しか利用されないため、利用率は低下します。

GPU は発散を測定するメトリックを提供し (通常、sub-group ごとにアクティブな work-item)、ネイティブの sub-group サイズと比較できます。

占有率

GPU の占有率については前述しましたが、これは簡単に言うと、特定のカーネルで実際にアクティブな sub-group の数と、アクティブな sub-group の理論上の最大数との比率です。占有率は、カーネルが利用可能な最大の並列性をどれくらい活用できるているかを開発者に示すことから重要です。

一部の GPU には、実際の占有率を測定するハードウェア機能が備わっています。理論上の占有率は、コンパイルされたカーネルとハードウェアのプロパティーから計算できます。

起動パラメーター

カーネルは、グローバルレンジとローカルレンジで起動されます。後者は work-group のサイズです。work-group のサイズは、sub-group サイズの倍数である必要があります。そのため、グローバル問題サイズを切り上げたり、グローバル問題サイズ外の work-item を処理しないようカーネルにコードを追加する必要があります。

占有率を改善するため、特定の GPU ハードウェアの work-group サイズに制約が課される場合があります。CU 数など、特定の GPU ハードウェアと何らかの関連性のあるグローバル問題サイズを選択することも有益な場合があります。グローバルとローカルの問題サイズは自然なサイズに合わせる必要がなく、ハードウェアに適合するように選択できます。

すべての GPU は、カーネルの起動ごとに実際の起動パラメーターを確認するメカニズムを提供しています。これには、グローバルおよびローカル問題サイズ、レジスター数、およびローカル・メモリー・サイズなどのカーネル・プロパティーが含まれます。

NVIDIA* GPU 上のパフォーマンス

このセクションでは、SYCL* プログラミング・モデルのコンテキスト内で NVIDIA* GPU アーキテクチャーと関連するパフォーマンスの考慮事項を説明します。NVIDIA* アーキテクチャー固有のパフォーマンスに関する最新の考慮事項については、NVIDIA* のドキュメント (英語) を参照してください。

アーキテクチャー

注意
ここでは、CUDA* と SYCL* 用語の簡単な対比を示しています。詳しい説明については、ComputeCpp の『SYCL* for CUDA* Developer』 (英語) を参照してください。ただし、ComputeCpp のガイドでは SYCL* 1.2 のみをカバーしているため、SYCL* 2020 で導入された USM など、いくつかの重要な機能については説明されていません。

NVIDIA* GPU の基本計算ユニットは、ストリーミング・マルチプロセッサーまたは SM と呼ばれます。SM は、32 個の work-item で構成される sub-group を実行します。NVIDIA では sub-group を ワープ (warp) と呼んでいます。work-item を実行するエンティティーはスレッド (thread) と呼ばれます。

work-group は、スレッドブロック (thread block) または協調スレッドアレイ (CTA: cooperative thread array) と呼ばれます。これには通常の規則が適用されます。CTA は同じ SM で同時に実行されることが保証され、SYCL* ローカルメモリー (CUDA* では共有メモリー) などのローカルリソースを利用でき、work-item 間で同期できます。

以下は、NVIDIA*/CUDA* と SYCL* の用語の対比表です。

NVIDIA*/CUDA* SYCL*
ストリーミング・マルチプロセッサー (SM) 計算ユニット (CU)
ワープ sub-group
スレッド work-item
スレッドブロック work-group
協調スレッドアレイ (CTA) work-group
共有メモリー ローカルメモリー
グローバルメモリー グローバルメモリー
ローカルメモリー プライベート・メモリー

 

work-group サイズは、sub-group サイズの 32 の倍数である必要があります。適切な work-group サイズは、占有率を最大化するように選択され、カーネルによって使用されるリソースに依存します。これは 1024 を越えることはありません。

NVIDIA* GPU では、デバイスクエリー sycl::info::device::sub_group_sizes が、値 32 の単一要素を持つベクトルを返します。

work-group サイズを選択したら、グローバルサイズを選択します。グローバルサイズを計算ユニット数の倍数にすると、負荷分散に役立つことがあります。デバイスクエリー sycl::info::device::max_compute_units は、計算ユニットの数を返します。

次に、それぞれの work-item (カーネル・インスタンス) が問題の複数の項目で動作するようにカーネルコードを記述します。起動パラメーターは、問題サイズにかかわりなく、ハードウェアのレイアウトに基づいて調整できます。

例えば、次の SYCL* コードは、ハードウェア・レイアウトに基づいて起動パラメーターを決定し、カーネル内のループで問題サイズを調整します。CUDA* では、このタイプの内部ループはグリッドストライド (grid-stride) と呼ばれます。


    int N = some_big_number;
    int wgsize = 256;
    int ncus = dev.get_info<info::device::max_compute_units>();
    int nglobal = 32 * ncus;
    cgh.parallel_for(nd_range<1>(nglobal * wgsize, wgsize),
      [=](nd_item<1> item)
      {
        int global_size = item.get_global_range()[0];
        for (int i = item.get_global_id(0); i < N; i += global_size)
          y[i] = a * x[i] + y[i];
      });

SM

NVIDIA* SM は、SM サブ・パーティション (英語) と呼ばれる 4 つの処理ブロックに分割されます。それぞれのワープは、存続期間全体で単一サブ・パーティションに存在します。ワープが停止すると、ハードウェア・ワープ・スケジューラーは準備ができている別のワープにスイッチできます。これは、いつでも実行できる十分な数のワープが必要であることを意味します。以降で、ハードウェア・メトリックを使用して、ワープ・スケジューラーの効率とストール時間を評価する方法を紹介します。

メモリー

メモリーにはいくつかの種類があります。

  • SYCL* グローバルメモリー (NVIDIA* グローバルメモリー) はデバイス上にありますが、一部のグローバルアドレスは、ホストまたは別のデバイス上にある CUDA* マネージドメモリー (SYCL* USM を使用) を参照する場合があります。すべてのグローバル・メモリー・アクセスは、すべての CU で共有される L2 キャッシュを経由します。カーネルの存続期間中のみ読み取られるデータは、sycl::ext::oneapi::experimental::cuda::ldg 関数を使用して CU ごとの L1 キャッシュに配置することができます。

  • SYCL* プライベート・メモリー (NVIDIA* ローカルメモリー) は、特定の work-item からのみアクセスできますが、グローバルメモリーにマップされているため、グローバルメモリーと比較してパフォーマンス上の優位性はありません。work-item ごとに同じアドレスにマップされます。これは、work-item スタック、レジスタースピル、およびその他の work-item のローカルデータに使用されます。

  • SYCL* ローカルメモリー (NVIDIA* 共有メモリー) は work-group で使用できます。グローバルメモリーよりも帯域幅が広く、レイテンシーが短縮されます。SYCL* ローカルメモリーには 32 のバンクがあります。連続する 32 ビット・ワードは、異なるバンクに割り当てられます。

  • SYCL* 共有メモリー (NVIDA* マネージドメモリー) は、ホストまたはアクセスされる任意のデバイス (上記のグローバルメモリーの説明を参照 ) に配置できます。ホストとデバイス間、またはデバイスとデバイス間でのメモリーの移行は、ランタイムによって管理されます。ただし、ユーザーは SYCL* 2020 API の prefetch や mem_advise を介して動作を制御するパフォーマンスのヒントを指定できます。

バンク競合

SYCL* ローカルメモリー内の 32 バイト float 配列にアクセスするストライド 1 のループを考えてみてください。最初の sub-group では、work-item 0 は配列要素 0 にアクセスし、work-item 1 は配列要素 1 にアクセスします。つまり、sub-group 全体として配列要素 0:31 にアクセスするため、バンク競合は起こりません。

一方、ループのストライドが 32 である場合、work-item は要素 32、64、96、… にアクセスし、sub-group は 32 要素離れた 32 の位置にアクセスするため、結果として 32 ウェイのバンク競合が発生します。これは、SYCL* ローカルメモリーの帯域幅を 1/32 に減少させます。

Volta のプレゼンテーション (英語) では、スライド 54 から 72 で共有メモリーバンクの競合について分かりやすく説明しています。

バンク競合を測定するハードウェア・メトリックがあります。

ローカル・メモリー・サイズ

デフォルトでは、カーネルはスレッドブロックあたり 48KB の SYCL* ローカルメモリーに制限されます。より大きな割り当てをサポートするプラットフォーム (Compute Capability 7.0 以降) では、SYCL_PI_CUDA_MAX_LOCAL_MEM_SIZE 環境変数を割り当て最大バイト数に設定して、明示的に有効にする必要があります。sycl::info::device::local_mem_size は、現在の SYCL* ローカルメモリーの最大値を返します。

Compute Capability 7.0 以降では、SYCL* ローカルメモリーと L1 キャッシュは、物理的に同じメモリーダイ上に配置されます。つまり、それらの合計は固定値であることを意味します。そのため、大量の共有メモリーを割り当てると、利用可能な L1 キャッシュが縮小してパフォーマンスに影響する可能性があります。

警告
Compute Capability 7.0 以降のカードでは、前述のように L1 キャッシュとローカルメモリーはハードウェアによって結合されています。共有メモリーはプラットフォーム機能によって任意のサイズに設定できますが、L1 にはそのような柔軟性はなく、固定値 (カーブアウトと呼ばれます) しか利用できません。そのため、ローカルメモリーを少し多く割り当てると、ドライバーは L1 を制限する次のカーブアウト値に切り替える可能性があります。

前述のように、メモリータイプ間の分割は実行時にドライバーによって行われます。sycl::local_accessor を使用して、割り当てられた実際のメモリー量に基づいて決定します。そのため、SYCL_PI_CUDA_MAX_LOCAL_MEM_SIZE を過度に大きく設定しても、選択した L1 キャッシュの量に影響することはありません。

特定の Compute Capability に対するこれらの適正値は、CUDA* C プログラミング・ガイド (英語) で確認できます。NVIDIA＊ では、「L1 キャッシュ」と「テクスチャー・キャッシュ」という用語が同じ意味で使用されることに注意が必要です。

メモリー結合

異なる work-item は通常、グローバルメモリー内の異なる位置にアクセスします。同じロード命令内の特定の sub-group でアクセスされるアドレスが、キャッシュラインの同じセットである場合、メモリーシステムは最小数のグローバルメモリーへのアクセスを発行します。これはメモリー結合と呼ばれます。この要件は、隣接するメモリー位置にアクセスする work-item によって容易に満たすことができます。データ構造を 32 バイト境界に配置することで、パフォーマンスをさらに向上できます。間接アクセス/大きなストライドによって、明らかにこれをが困難であることがあります。

最近の NVIDIA* GPU では、洗練されたキャッシュとメモリーシステムを備えているため、グローバル・メモリー・アクセスが結合される可能性は高くなります。結合の測定に有効なハードウェア・メトリックがあります。

キャッシュ

すべての GPU ユニットは L2 キャッシュを共有します。L2 キャッシュは物理アドレスでアクセスされます。これには、データ圧縮とグローバルアトミック (浮動小数点加算など) の機能も含まれます。

各 SM には、複数の機能に使用される L1 キャッシュがあります。L1 のスループットは、パフォーマンスを制限する要因になることがあります。

占有率

NVIDIA* では、同時にアクティブな CU 数を、利用可能な最大 CU 数で割ったものを占有率として定義しています。明らかに、占有率が上がると GPU 利用率が上がり、パフォーマンスが向上することが期待されます。理論上の占有率は、ハードウェアの制限、カーネルで使用されるレジスター数、およびカーネルで使用される共有メモリーの量によって決定されます。

理論上の占有率は、work-group サイズと特定のカーネルを起動するパラメーターを決定するガイドラインとなります。NVIDIA* プロファイリング・ツールは、起動されるカーネルごとの実際の占有率と理論上の占有率を示します。実行時に work-group サイズを変更して、結果をベンチマークして最適な work-group サイズを選択できます。しかし、コードを実行することなく理論的な占有率の推測を利用することもできます。

NVIDIA は、オンラインの spreadsheet を提供しています。公式には非推奨ですが、NVIDIA* Nsight* Compute の占有率セクションと同じ機能を提供し、理論上の占有率がどのように決定されるかを理解するのに役立ちます。代わりに NVIDIA* Nsight* Compute (英語) を使用することが推奨されています。

パフォーマンス・ツール

一般に、NVIDIA* が CUDA* 向けに提供するすべてのパフォーマンス・ツールは、DPC++ CUDA* プラグインを使用して SYCL* アプリケーションをシームレスに処理します。

この節では、それらのパフォーマンス・ツールの一部を紹介します。

NVIDIA Nsight* Systems (nsys)

NVIDIA Nsight* Systems (英語) は、コマンドライン・ツール nsys と GUI nsys-ui を含んでいます。さらに、NVTX (英語) トレース・ライブラリーをアプリケーションで使用して、nsys 解析の対象領域を制限し、アプリケーション固有のイベントを nsys 解析に追加することができます。

基本的な使い方

$ nsys profile <command> <arguments>

このコマンドは、指定するコマンドを与える引数で実行し、デフォルト設定でプロファイルを行います。出力は、NVIDIA* 固有の形式でレポートファイルに書き込まれます。

nsys はシステム全体を監視することに注意してください。コマンドは何であっても構いませんが、nsys は単にコマンドを起動して監視を開始し、コマンドが終了するまで監視を続行します。

$ nsys stats <report file>

このコマンドは、レポートファイルから sqlite データベースを作成し、データベースでいくつかのレポートを作成します。レポートには、API の使用状況、GPU に関連するメモリーコピー、および GPU カーネルのタイミングなどが含まれます。結果はデフォルトでコンソールに出力されます。レポートを csv 形式のファイルに書き込むオプションもあります。

nsys-ui コマンドを使用して、レポートファイルを GUI で調査することもできます。これには、ローカルまたは VNC のようなリモート X ビューアーを介して接続されたディスプレイが必要です。レポートファイルは、ラップトップなどのローカルマシンに移動して、そこで GUI を実行できます。レポートファイルは非常に大きくなる可能性があり、大量のメモリーを必要とする場合があります。

アプリケーションで NVTX アノテーションを戦略的に使用し、コマンドライン引数を追加することで、レポートサイズを縮小し、解析ワークフローを円滑にすることができます。

レポートは JSON 形式でエクスポートすることもできます。これには、sqlite データベースと同じデータが含まれていますが、SQL を熟知していない開発者にも解析が容易です。

NVTX アノテーション

NVTX は、NVIDIA* が提供するインストルメント API で、パフォーマンス・ツールで利用されます。この API は、nvToolsExt.h をインクルードすることで利用できます。ライブラリーのリンクなどは必要ありません。

NVTX は、大規模なアプリケーションで、ほかの NVIDIA* ツールのデータ収集をアプリケーションの特定の部分に制限するのに役立ちます。

NVTX は、デバイスコードではなく、ホストコードの測定にのみ利用できることに注意してください。

NVTX の使い方と機能の詳細については、NVIDIA* のドキュメント (英語) を参照してください。

NVIDIA* Nsight* Compute (ncu)

前の節では、NVIDIA Nsight* Systems (nsys) について説明しました。NVIDIA* Nsight* Compute (ncu) は、GPU ハードウェアのパフォーマンスに注目する支援ツールです。ncu を使用すると、NVIDIA* GPU で利用可能なハードウェア・カウンターにアクセスできるだけでなく、特定のカーネルが GPU をどの程度活用しているかを理解するのに役立つ事前定義された解析タイプ (セクションと呼ばれます) を知ることができます。

ncu には GUI 版がありますが、ここでは nsys の説明のようにコマンドラインといくつかの処理スクリプトを利用します。ncu CLI については、https://docs.nvidia.com/nsight-compute/NsightComputeCli/index.html (英語) で説明されています。

ncu オーバーヘッド

ncu は、アプリケーションの実行時間を大幅に増加させる可能性があります。いくつかの原因が考えられます。

  • 指定されたメトリックを収集するため、場合によってはカーネルを複数回実行する必要があります。

  • 一部のメトリックでは、オーバーヘッドが高いカーネルのバイナリー・インストルメンテーションが必要になります。

  • 通常、カーネルの実行はシリアル化されるため、同時実行性が低下します。

詳細については、「カーネル・プロファイル・ガイド」 (英語) を参照してください。

速度が低下すると、通常、収集プロセスを制限する必要があります。NVIDIA* では、これに対処するいくつかの方法を提案しています。

  • nsys の節で説明したように、NVTX を使用してアプリケーションをインストルメントし、ncu にコマンドライン・オプション --nvtx および --nvtx-include または --nvtx-exclude を指定して、特定の NVTX レンジを含めるか除外するかを指示します。ドメイン内のレンジを指定する構文は、nsys とは逆であることに注意してください。range@domain ではなく domain@range となります。ncu の構文は非常に豊富ですが複雑です。「NVTX フィルター処理」 (英語) のセクションを参照してください。

  • -k kernelname を指定して、単一カーネルのデータのみを収集します。カーネル名には正規表現を利用できます。
    注意: C++ ユーザーは、--kernel-name-base=mangled を指定してマングルされた名前を使用することもできます。

  • --launch-count と --launch-skip を使用して、特定数のカーネルの起動を収集します。

上記は組み合わせて使用できます。

セクション

ncu には、セクションと呼ばれる事前定義された多数のメトリックのセットが用意されています。各セクションは、特定のパフォーマンスに関する疑問を解決するのを支援します (例: アプリケーションはメモリー依存であるか、など)。

セクションは、ncu --list-sections で一覧を表示できます。ncu バージョン 2022.1.1.0 の 出力を以下に示します。

識別子 表記名
ComputeWorkloadAnalysis Compute Workload Analysis
InstructionStats Instruction Statistics
LaunchStats Launch Statistics
MemoryWorkloadAnalysis Memory Workload Analysis
MemoryWorkloadAnalysis_Chart Memory Workload Analysis Chart
MemoryWorkloadAnalysis_Deprecated (Deprecated) Memory Workload Analysis
MemoryWorkloadAnalysis_Tables Memory Workload Analysis Tables
Nvlink NVLink
Nvlink_Tables NVLink Tables
Nvlink_Topology NVLink Topology
Occupancy Occupancy
SchedulerStats Scheduler Statistics
SourceCounters Source Counters
SpeedOfLight GPU Speed Of Light Throughput
SpeedOfLight_HierarchicalDoubleRooflineChart GPU Speed Of Light Hierarchical Roofline Chart (Double Precision)
SpeedOfLight_HierarchicalHalfRooflineChart GPU Speed Of Light Hierarchical Roofline Chart (Half Precision)
SpeedOfLight_HierarchicalSingleRooflineChart GPU Speed Of Light Hierarchical Roofline Chart (Single Precision)
SpeedOfLight_HierarchicalTensorRooflineChart GPU Speed Of Light Hierarchical Roofline Chart (Tensor Core)
SpeedOfLight_RooflineChart GPU Speed Of Light Roofline Chart
WarpStateStats Warp State Statistics

メトリック

セクションの代わりに、--metrics を使用して特定のメトリックを収集するよう ncu に指示できます。利用可能なメトリックは、--query-metrics で照会できます。

出力

この節の例では、解析が容易な出力を生成するため、--csv オプションを使用しています。--log-file オプションは、この出力と他の出力をともに stdout の代わりにファイルへ送ります。

プロファイル・レポートは、ncu GUI によって使用されます。これは、--export オプションで保存できます。その後、ファイルを別のマシンに移動して ncu-gui を使用し、ローカルで表示できます。

分断された名前は、--print-kernel-base=mangled で選択できます。

カーネル・アセンブリーの抽出

状況によっては、特定のカーネルのパフォーマンスを理解するため、アセンブリーを調べることが有効な場合があります。NVIDIA* の場合は、コンパイラーがカーネル向けに生成した PTX を調査します。

NVIDIA* GPU 用にビルドされた SYCL* アプリケーションから PTX を抽出するには、環境変数 SYCL_DUMP_IMAGES に 1 を設定してアプリケーションを実行します。これにより、現在の作業ディレクトリーに sycl_nvptx641.bin のような名前のファイルが生成されます。これらのファイルは CUDA* fat バイナリーであり、NVIDIA* ターゲットに依存しない仮想アセンブリー言語である PTX と、単一または複数のターゲットのマシンコードである SASS が含まれます。

次のように CUDA* ツール cuobjdump (英語) を使用して、FAT バイナリーから PTX と SASS の両方を抽出できます。

# Extract PTX
cuobjdump --dump-ptx sycl_nvptx641.bin

# Extract SASS
cuobjdump -sass sycl_nvptx641.bin

NVIDIA* Nsight* Compute ツールの GUI バ―ションでも逆アセンブリーを表示できます。

一般的な最適化

ここでは、DPC++ を使用する際の一般的なパフォーマンスの問題や落とし穴、そしてその対処方法について説明します。

インデックスの入れ替え

SYCL* 仕様の 4.9.1 節では、次のことが規定されています。

整数から多次元 id やレンジを構成する場合、多次元空間の線形化において右端の要素が最も速く変化するように要素を記述します。

そのため、インテル® DPC++ コンパイラーでは、右端の次元が CUDA* または HIP の x 次元にマップされ、右から 2 つ目の次元が CUDA* や HIP の y 次元にマップされます。以下に例を示します。

cgh.parallel_for(sycl::nd_range{sycl::range(WG_X), sycl::range(WI_X)}, ...)

cgh.parallel_for(sycl::nd_range<2>{sycl::range<2>(WG_Y, WG_X), sycl::range<2>(WI_Y, WI_X)}, ...)

cgh.parallel_for(sycl::nd_range<3>{sycl::range<3>(WG_Z, WG_Y, WG_X), sycl::range<3>(WI_Z, WI_Y, WI_X)}, ...)

WG_X と WI_X は、x 次元の work-group 数と work-group ごとの work-item 数 (CUDA* では、名前付きのグリッドサイズとブロックあたりのスレッド数) であり、_Y と _Z は y 次元と z 次元のものになります。

次の場合、2 次元または 3 次元のカーネルの parallel_for 実行では特に重要であることに注意してください。

  • ローカルまたはグローバルメモリー内の (1-d) 配列には、手動で線形化されたアクセスがあります。結合されていないグローバル・メモリー・アクセスまたはローカルメモリー内のバンク競合によるパフォーマンスの問題を回避するため、これを考慮する必要があります。線形化の詳細については、SYCL* 仕様の 3.11 節の多次元オブジェクトと線形化を参照してください。

  • 次のエラー (または同等のエラー) が発生します。

    Number of work-groups exceed limit for dimension 1 : 379957 > 65535

    これは、CUDA* など一部のプラットフォームでは、x 次元が y および z 次元よりも多くの work-group をサポートするためです。

    Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)

このトピックの詳細については、こちら (英語) を参照してください。

特定の GPU アーキテクチャーの設定

NVIDIA* GPU 向けにビルドする場合、GPU のアーキテクチャーは省略できます。この場合、デフォルトのアーキテクチャーは sm_50 になります。新しい世代の GPU を使用する場合、適切なアーキテクチャーを指定することでコンパイラーがより高性能な新しい機能を使用できる可能性があります。特定の GPU アーキテクチャーを指定する方法は、「導入ガイド」を参照してください。

インライン展開

DPC++ は、パフォーマンスとコンパイル時間のバランスを考慮して、自動的に関数をインライン展開するかどうかを選択します。しかし、プログラマーは、特定の関数に always_inline 属性を追加するなどして、インライン展開を強制することもできます。

__attribute__((always_inline)) void function(...) {
  ...
}

...

q.submit([&](sycl::handler &cgh) {
    cgh.parallel_for(..., [=](...) {
      function(...);
    });
}

高速数学ビルトイン

SYCL* 数学ビルトインは、同等の OpenCL* 1.2 数学ビルトインの精度要件と一致するように定義されていますが、一部のアプリケーションでは必要以上に精度が高くなり、パフォーマンスが低下する可能性があります。

これに対処するため、SYCL* 仕様では、数学関数のサブセットのネイティブバージョン (4.17.5 節の「数学関数」に完全なリストがあります) が提供されています。これには、精度とパフォーマンスのトレードオフがあります。これらは、ネイティブ名前空間で定義されています。例えば、sycl::cos() のネイティブバージョンは、sycl::native::cos() です。

一般に、精度が問題にならない場合、ネイティブバリアントを使用すると大幅に改善できる可能性がありますが、すべてのバックエンドがすべてのビルトインに対し緩和された精度を使用するわけではないことに注意してください。

NVIDIA* アーキテクチャーをターゲットにする場合、次の sycl::native:: 関数は、対応する PTX 命令の .approx バリアントによって実装されます。

  • sycl::native::divide

  • sycl::native::sqrt

  • sycl::native::sin

  • sycl::native::cos

  • sycl::native::log2

  • sycl::ext::oneapi::experimental::native::exp2

  • sycl::ext::oneapi::experimental::native::tanh

上記 sycl::native:: 数学関数と PTX 命令は 1:1 でマッピングされます。例えば、1 つの sycl::native::exp2 呼び出しに対し、コンパイラーは 1 つの ex2.approx 命令を生成します。別のケースでは、ネイティブ数学関数は、複数の .approx PTX 命令によって実装されます。例えば、sycl::native::tan() は、sin.approx、cos.approx、および divide.approx で実装されます。ネイティブ数学関数は、精度と引き換えに sycl:: の対応する関数よりも高速になる可能性があります。.approx PTX 命令の精度の詳細については、PTX ISA ドキュメント (英語) を参照してください。

注意
-ffast-math コンパイルオプションは、標準の sycl:: 数学関数を、対応する sycl::native:: 関数に入れ替えます (利用可能であれば)。指定された数学関数のネイティブバージョンが存在しない場合、-ffast-math フラグは影響しません。

icpx コンパイラーでは -ffast-math がデフォルトです。icpx で -ffast-math を無効にするには、-fno-fast-math を使用します。

ループアンロール

コンパイラーは一部のループアンロールを自動的に行いますが、次のように unrolling プラグマを使用して、デバイスコード内の計算集約型ループのアンロールを手動でコンパイラーに指示することが有益な場合もあります。

#pragma unroll <unroll factor>
for( ... ) {
 ...
}

エイリアス解析

エイリアス解析では、2 つのメモリー参照が互いにエイリアスでないことが証明できます。これにより最適化が有効になることがあります。デフォルトでは、コンパイラーはエイリアス解析によって証明されない限り、メモリー参照はエイリアスであると想定する必要があります。ただし、デバイスコード内のメモリー参照がエイリアスではないことをコンパイラーに明示的に通知することもできます。これは、バッファー/アクセサーと USM モデルのそれぞれのキーワードを使用することで実現できます。

前者は、oneapi 拡張の no_alias プロパティーをアクセサーに追加することができます。

q.submit([&](sycl::handler &cgh) {
  sycl::accessor acc{..., sycl::ext::oneapi::accessor_property_list{sycl::ext::oneapi::no_alias}};
  ...
});

後者の場合、__restrict__ 修飾子をポインターに追加できます。

__restrict__ は C++ では非標準であり、SYCL* 実装全体で一貫性がない可能性があることに注意してください。dpc++ では、restrict 修飾されたデバイス関数 (SYCL* カーネルから呼び出される関数) パラメーターのみが考慮されます。

例:

void function(int *__restrict__ ptr) {
    ...
}

...
int *ptr = sycl::malloc_device<int>(..., q);
...
q.submit([&](sycl::handler &cgh) {
      cgh.parallel_for(..., [=](...) {
        function(ptr);
      });
});

より強制的なアプローチは、[[intel::kernel_args_restrict]] 属性をカーネルに追加することです。これは、各 USM ポインター間、またはそのモデルがカーネル内で使用される場合はバッファーアクセサー間のすべてのエイリアス依存関係を無視するようにコンパイラーに指示します。

例 (バッファー/アクセサーモデル):

q.submit([&](handler& cgh) {
      accessor in_accessor(in_buf, cgh, read_only);
      accessor out_accessor(out_buf, cgh, write_only);
        cgh.single_task<NoAliases>([=]() [[intel::kernel_args_restrict]] {
          for (int i = 0; i < N; i++)
              out_accessor[i] = in_accessor[i];
        });
});

テクスチャー・キャッシュの使用

CUDA* プラットフォームでは、少なくともカーネルの存続期間中は一定であるデータをテクスチャー・キャッシュにキャッシュできます。

これは、sycl::ext::oneapi::experimental::cuda::ldg 関数を使用して実現できます。この関数は、デバイスメモリーへのポインターを受け取り、L1/tex キャッシュからロードして、アドレスに格納された値を返します。以下に例を示します。

float some_value = ldg(&some_data_in_device_memory[some_index]);

警告
この関数でロードされたデータがカーネル内に書き込まれることをコンパイラーが検出した場合でも、プログラムはコンパイルできますが、テクスチャー・キャッシュは使用されないことに注意してください。

テクスチャー・キャッシュを使用することでパフォーマンス向上に影響する要因は数多くあります。そのため、最大の高速化を達成するのは困難な場合があります。実際、多くのユースケースではメリットがほとんどないか、全くありません。ただし、パフォーマンスが低下する可能性は低く、低下した場合でも規模は小さく、ldg は最小限のコード変更で使用できるため、カーネルのパフォーマンスを素早く向上する素晴らしい方法となるかもしれません。

ldg は、HIP AMD を含むほかのすべてのプラットフォームでも移植可能であることに注意してください。しかし、CUDA* は現在 ldg により特殊なキャッシングが可能な唯一のプラットフォームです。HIP AMD バックエンドは、ldg を使用するかどうかにかかわらず、常にすべてのレジスターデータを L1 キャッシュと L2 キャッシュにロードします。

テクスチャー・キャッシュの詳細については、こちらのブログ (英語) を参照してください。ldg 関数の詳細は、対応する拡張機能のドキュメント (英語) をご覧ください。

 

サポート

機能

コア機能

機能 サポート
コンテキスト内の複数デバイス いいえ
サブグループ (sub-group) 部分的[1]
グループ関数/アルゴリズム はい
整数関数 はい
数学関数 (スカラー) はい
数学関数 (ベクトル) はい
数学関数 (marray) いいえ
共通関数 はい
ジオメトリー関数 はい
リレーショナル関数 はい
atomic ref はい
オペレーティング・システム Linux*
バッファーの再解釈 はい
stream はい
デバイスイベント はい
グループの非同期コピー はい
プラットフォームの get info はい
カーネルの get info はい
sycl::nan と sycl::isnan はい
デバイスセレクター はい
階層的並列化 はい
ホストタスク はい
インオーダー・キュー はい
リダクション はい
キューのショートカット はい
vec はい
marray はい
errc はい
匿名カーネルラムダ はい
機能を評価するマクロ はい
sycl::span はい
sycl::dynamic_extent いいえ[2]
sycl::bit_cast はい
aspect_selector いいえ
カーネルバンドル いいえ
特殊化定数 はい

非コア機能

機能 サポート
image いいえ (SYCL_PI_CUDA_ENABLE_IMAGE_SUPPORT の一部)
fp16 データタイプ はい
fp64 データタイプ はい
prefetch はい
USM ホスト、デバイス、共有
USM アトミックホスト割り当て いいえ
USM アトミック共有割り当て はい
USM システムに割り当て はい
SYCL_EXTERNAL はい
アトミックメモリーの順序付け relaxed、acquire、release、acq_rel
アトミック・フェンス・メモリーの順序付け いいえ
アトミック・メモリー・スコープ sub_group、work_group、device、system
アトミック・フェンス・メモリーのスコープ いいえ
64 ビット・アトミック はい
バイナリー形式 NVPTX と SASS
デバイスのパーティション化 いいえ
ホストデバッグ可能デバイス いいえ
オンラインコンパイル いいえ
オンラインリンカー はい
キューのプロファイル はい
mem_advise read_mostly、優先場所、アクセス元
バックエンド仕様 WIP[3]
アプリケーション・バックエンドの相互運用 部分的[3]
カーネル・バックエンドの相互運用 いいえ
ホストタスク (ハンドルと相互運用) はい
reqd_work_group_size 部分的[4]
キャッシュビルド結果 いいえ
ビルドログ いいえ
ビルトインカーネル関数 なし

拡張機能

機能 サポート
uniform いいえ
USM アドレス空間 (デバイス、ホスト) 部分的[5]
固定ホストメモリーの使用 はい
サブグループ・マスク (+ グループ投票) はい
静的ローカルメモリー使用量照会 いいえ
sRGB イメージ いいえ
デフォルト・プラットフォーム・コンテキスト 一部分
メモリーチャネル いいえ
最大ワークグループ照会 一部分
結合行列 完全 [^jointmatrix、experiemtal]
すべて制限 (restrict all) いいえ
プロパティー・リスト (property list) いいえ
カーネル・プロパティー (kernel properties) いいえ
SIMD 呼び出し いいえ
低レベルデバイス情報 いいえ
カーネルキャッシュ設定 いいえ
FPGA lsu いいえ
FPGA reg いいえ
データ・フロー・パイプ いいえ
キューに投入されたバリア いいえ
フィルターセレクター はい
グループソート はい
フリー関数の照会 はい
明示的な SIMD いいえ
discard_queue_events 部分的 [1]
device_if いいえ
device_global いいえ
C と C++ 標準ライブラリーのサポート はい (clang++ 使用時)
カーネルでの assert はい
buffer_location いいえ
accessor_property_list (+ no_offset、no_alias) はい
グループ・ローカル・メモリー はい
printf はい
ext_oneapi_bfloat16 はい
拡張デバイス情報 いいえ
sycl_ext_oneapi_cuda_tex_cache_read はい
sycl_ext_oneapi_native_math はい
sycl_ext_oneapi_bfloat16_math_functions はい
sycl_ext_oneapi_cuda_async_barrier はい
sycl_ext_oneapi_bindless_images いいえ
sycl_ext_oneapi_graph はい
sycl_ext_oneapi_non_uniform_groups はい (tangle_group を除く)
sycl_ext_oneapi_peer_access はい

 

[1] (1、2、3) 一部のテストで失敗
[2] numeric_limits<size_t>::max() の使用
[3] (1、2) https://github.com/KhronosGroup/SYCL-Docs/pull/197 (英語) を参照
[4] https://github.com/intel/llvm/issues/6103 (英語) を参照
[5] https://github.com/intel/llvm/pull/6289 (英語) に追加 (未テスト)

 

更新履歴

2024.0

改良点

SYCL* コンパイラー

  • sycl::range を使用する parallel_for インターフェイスの work_group サイズの選択が改善されました。

SYCL* ライブラリー

  • sycl_ext_oneapi_graph 拡張がサポートされました。

  • sycl_ext_oneapi_non_uniform_groups 拡張がサポートされました。

  • sycl_ext_oneapi_peer_access 拡張がサポートされました。

  • max_registers_per_work_group デバイスクエリーが導入されました。

  • SYCL_PROGRAM_COMPILE_OPTIONS を使用するメカニズムが追加され、maxrregcount ptxas コンパイラー・オプションを次のように CUDA* バックエンドに渡すことができるようになりました。

    SYCL_PROGRAM_COMPILE_OPTIONS="--maxrregcount=<value>"

    これは、JIT コンパイルされたプログラムに対してのみ機能することに注意してください。

バグフィックス

  • 倍精度浮動小数点グループ・アルゴリズムが修正され、icpx コンパイラーでコンパイルしてもハングアップしなくなりました。
    broadcast、joint_exclusive_scan、joint_inclusive_scan、exclusive_scan_over_group、inclusive_scan_over_group

  • 引数として sycl::memory_scope::device を渡したときに、誤った memory_scope を使用していた atomic_fence のバグを修正しました。

2023.2.0

改良点

SYCL* コンパイラー

  • NVPTX バックエンドのインライン展開のしきい値乗数を増やすことで clang++/cuda のパフォーマンスが向上しました [22d98280]。

  • SYCL* の __CUDA_ARCH__ の代わりに __SYCL_CUDA_ARCH__ を定義しました [8f5000c3]。

SYCL* ライブラリー

  • __ldg* clang ビルトインを CUDA* 専用の拡張機能として SYCL* に公開する sycl_ext_oneapi_cuda_tex_cache_read を導入しました – 読み取り専用テクスチャー・キャッシュ (英語) [5360825e]。

  • cl_khr_subgroups 拡張機能をサポートするサブグループとしてレポートするようになりました [8e6c092b]。

  • atomic_fence デバイスクエリーは、NVIDIA によってドロップされる可能性があるエラー [82ac98f8] で失敗せずに、必要最小限の機能を返すようになりました [1e88df54]。

  • 理論上のピークメモリー帯域幅のクエリーをサポートしました – インテルのデバイス情報拡張機能 (英語) [8ce0a6d5]

  • デバイス ID と UUID のサポートを追加しました – インテルのデバイス情報拡張機能 (英語) [8213074d]

  • ホストデバイス memcpy2D をサポートしました [d0b25d4a]。

  • CUDA* バックエンドで sycl_ext_oneapi_memcpy2d をサポートしました – oneAPI memcpy2d (英語) [9008a5d2]

バグフィックス

  • 無効な work-group サイズに関するエラーを PI_ERROR_INVALID_WORK_GROUP_SIZE に置き換えるようになりました [2357af0a]。

  • sycl::ctz 関数からの間違った結果に対応しました [5a9f601e]。

  • イベントが意図したとおりに待機しない原因となる問題に対処しました [1b225447]。

2023.1.0

改良点

SYCL* コンパイラー

  • FTZ、prec-sqrt が no-ftz、no-prec-sqrt をオーバーライドできるようになりました [8096a6fb]。

  • -fsycl-targets の引数として NVIDIA* アーキテクチャー (nvidia_gpu_sm_80 など) を指定できるようになりました [e5de913f]。

SYCL* ライブラリー

  • 新しい「unified」インターフェイスによる行列拡張を実装しました [166bbc36]。

  • CUDA* バックエンドのゼロ・レンジ・カーネルをサポートしました [a3958865]。

  • 不足しているマクロを interop-backend-traits.cpp に追加しました [a578c8141]。

  • CUDA* バックエンドで各種メタデータを許可するようになりました [25d05f3d]。

  • tf32 デバイスコードのチェックコメントを更新しました [21176576]。

バグフィックス

  • ext_oneapi_cuda make_device が sycl::device を複製しなくなりました [75302c53a]。

  • ガードが正しく構築されない問題を修正しました [ce7c594f]。

ドキュメント

  • ptxas オプションを渡す方法がドキュメント化されました [f48f96eb3f]。

  • cuda-arch 固有の機能を有効にする CUDA* GPU アーキテクチャーの説明が追加されました [4e5d276f]。

2023.0.0

oneAPI for NVIDIA* GPU の最初のリリースです。

このリリースは、intel/llvm repository at commit 0f579ba (英語) から作成されました。

新機能

  • CUDA* バックエンドのサポート

SYCL* コンパイラー

  • sycl::half タイプのサポート
  • ストレージタイプで動作する bf16 ビルトインのサポート
  • リレーショナル、ジオメトリー、共通、および数学カテゴリーからの SYCL* ビルトインのサポート
  • sub_group 拡張のサポート
  • グループ・アルゴリズムのサポート
  • group_ballot 組込み関数のサポート
  • スコープとメモリー順序によるアトミックのサポート
  • 同時実行を改善する各キューでの複数ストリームのサポート
  • sycl::queue::mem_advise のサポート
  • CUDA* libclc での --ffast-math のサポート
  • デバイスでの assert をサポート
  • CUDA* libclc での float/double 変換、比較交換のアトミック操作をサポート
  • CXX 標準ライブラリー関数に対応
  • デフォルト ctor の sycl::event のネイティブイベントは COMPLETE 状態

SYCL* ライブラリー

  • fma、fmin、fmax および fmax に bf16 ビルトインを追加
  • sycl::aspect::fp16 をサポート
  • tanh (float/half) と exp2 (half) のネイティブ定義を追加
  • sycl::get_native(sycl::buffer) をサポート
  • mem_advise リセットと同時メモリーチェック管理を実装
  • bfloat16 のサポートを含む、joint_matrix での要素ごとの操作をサポート
  • 統合共有メモリー (USM) をサポート

 

トラブルシューティング

この節では、トラブルシューティングのヒントと一般的な問題の解決方法について説明します。ここで説明する方法で問題が解決しない場合は、Codeplay のコミュニティー・サポート・ウェブサイト (英語) からサポートリクエストをお送りください。完全なサポートは保証できませんが、できる限り支援させていただきます。サポートリクエストを送信する前に、ソフトウェアが最新の安定したバージョンであることを確認してください。

問題、パフォーマンス、機能要望は、oneAPI DPC++ コンパイラーのオープンソース・リポジトリー (英語) からも報告できます。

sycl-ls の出力にデバイスが表示されない

sycl-ls がシステム上の期待されるデバイスを報告しない場合:

  1. システムに互換性のあるバージョンの CUDA* または ROCm* ツールキット (それぞれ CUDA* と HIP プラグイン向け)、および互換性のあるドライバーがインストールされていることを確認してください。

  2. nvidia-smi または rocm-smi がデバイスを正しく認識できることを確認します。

  3. プラグインが正しくロードされていることを確認します。これは、環境変数 SYCL_PI_TRACE に 1 を設定して、sycl-ls を再度実行することで分かります。

    例:

    $ SYCL_PI_TRACE=1 sycl-ls

    次のような出力が得られるはずです。

    SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_opencl.so [ PluginVersion: 11.15.1 ]
SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_level_zero.so [ PluginVersion: 11.15.1 ]
SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_cuda.so [ PluginVersion: 11.15.1 ]
[ext_oneapi_cuda:gpu:0] NVIDIA CUDA BACKEND, NVIDIA A100-PCIE-40GB 0.0 [CUDA 11.7]

    インストールしたプラグインが sycl-ls の出力に表示されない場合、SYCL_PI_TRACE に -1 を設定して再度実行することで、詳細なエラー情報を取得できます。

    $ SYCL_PI_TRACE=-1 sycl-ls

    大量の出力が得られますが、次のようなエラーが表示されているか確認してください。

    SYCL_PI_TRACE[-1]: dlopen(/opt/intel/oneapi/compiler/2024.0.0/linux/lib/libpi_hip.so) failed with <libamdhip64.so.4: cannot open shared object file: No such file or directory>
SYCL_PI_TRACE[all]: Check if plugin is present. Failed to load plugin: libpi_hip.so

    • CUDA* プラグインには、CUDA* SDK で提供される libcuda.so と libcupti.so が必要です。

    • HIP プラグインには、ROCm* の libamdhip64.so が必要です。

    CUDA* または ROCm* のインストールと、環境が適切に設定されていることを確認してください。また、LD_LIBRARY_PATH が上記のライブラリーを検出できる場所を指しているか確認してください。

  4. ONEAPI_DEVICE_SELECTOR または SYCL_DEVICE_ALLOWLIST などのデバイスフィルター環境変数が設定されていないことを確認します (ONEAPI_DEVICE_SELECTOR が設定されていると、sycl-ls は警告を表示します)。

  5. 権限を確認します。POSIX* では、アクセラレーター・デバイスへのアクセスは、通常、適切なグループのメンバーであることを条件としています。例えば、Ubuntu* Linux* の場合、GPUへのアクセスには video グループと render グループのメンバーである必要がありますが、これは設定によって異なります。

 

不正バイナリーエラーの扱い

不適切なプラットフォーム

よくある間違いは、SYCL* プログラムに互換性のあるバイナリーがないプラットフォームを使用して SYCL* プログラムを実行することです。例えば、SYCL* プログラムは SPIR-V* バックエンド用にコンパイルされた後、HIP デバイス上で実行される可能性があります。この場合、PI_RESULT_ERROR_INVALID_BINARY エラーコードがスローされます。この場合、次の点を確認してください。

  1. プログラムが適切なプラットフォーム向けにコンパイルされるよう、-fsycl-targets にターゲット・プラットフォームが指定されていることを確認します。

  2. プログラムが、実行可能ファイルがコンパイルされたプラットフォームと互換性のある sycl プラットフォームまたはデバイスセレクターを使用していることを確認します。

適切なプラットフォームと不適切なデバイス

CUDA* または HIP をターゲットにする SYCL* アプリケーションを実行すると、特定の状況でアプリケーションが失敗し、無効なバイナリーであることを示すエラーが報告されることがあります。例えば、CUDA* の場合は CUDA_ERROR_NO_BINARY_FOR_GPU がレポートされる場合があります。

これは、選択された SYCL* デバイスに適切でないアーキテクチャーのバイナリーが送信されたことを意味します。この場合、次の点を確認してください。

  1. アプリケーションが、利用するハードウェアのアーキテクチャーと一致するようにビルドされていることを確認してください。

    • CUDA* 向けのフラグ: -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=<arch>

    • HIP 向けのフラグ: -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=<arch>

  2. 実行時に適切な SYCL* デバイス (ビルドされたアプリケーションのアーキテクチャーに一致するもの) が選択されていることを確認します。環境変数 SYCL_PI_TRACE=1 を設定すると、選択されたデバイスに関連するトレース情報を表示できます。以下に例を示します。

    SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_opencl.so [ PluginVersion: 11.16.1 ]
SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_level_zero.so [ PluginVersion: 11.16.1 ]
SYCL_PI_TRACE[basic]: Plugin found and successfully loaded: libpi_cuda.so [ PluginVersion: 11.16.1 ]
SYCL_PI_TRACE[all]: Requested device_type: info::device_type::automatic
SYCL_PI_TRACE[all]: Requested device_type: info::device_type::automatic
SYCL_PI_TRACE[all]: Selected device: -> final score = 1500
SYCL_PI_TRACE[all]:   platform: NVIDIA CUDA BACKEND
SYCL_PI_TRACE[all]:   device: NVIDIA GeForce GTX 1050 Ti

  3. 誤ったデバイスが選択されている場合、環境変数 ONEAPI_DEVICE_SELECTOR を使用して SYCL* デバイスセレクターが選択するデバイスを変更できます。インテル® oneAPI DPC++/C++ コンパイラーのドキュメントにある「環境変数」の節を参照してください。

外部参照関数「…」を解決できません/外部シンボル「…」が未定義です

これにはいくつかの原因が考えられます。

  1. 現在 DPC++ では std::complex はサポートされていません。代わりに sycl::complex を使用してください。

  2. icpx コンパイラーは、デフォルトで -ffast-math モードを使用するため、現在 ldexp や logf などの特定の数学関数の解決に問題が生じることがあります。これは、-fno-fast-math フラグを使用して -ffast-math を無効にすることで回避できます。

    詳細は、「oneAPI for NVIDIA* GPU のインストール」を参照してください。

コンパイラーのエラー: “cannot find libdevice (libdevice が見つかりません)”

CUDA* SDK がデフォルトの位置にインストールされていないと、clang++ が SDK を検出できず、コンパイル中に次のようなエラーが発生することがあります。

clang-17: error: cannot find libdevice for sm_50; provide path to different CUDA installation via '--cuda-path', or pass '-nocudalib' to build without linking with libdevice

この問題を解決するには、--cuda-path オプションで CUDA* SDK のインストール・パスを指定します。

コンパイラーのエラー: “needs target feature (ターゲットの機能が必要です)”

DPC++ ランタイムで使用される一部の nvptx ビルトインは、コンパイルに最小限の計算機能を要求します。プログラムが使用するビルトインに対し十分な計算機能をターゲットにしていない場合 (コンパイル引数 -Xsycl-target-backend --cuda-gpu-arch=sm_xx を使用)、次のエラーが報告されます。

error: '__builtin_name' needs target feature (sm_70|sm_72|..),...

このようなエラーを回避するには、十分な計算機能を持つデバイスをターゲットにしてコンパイルしていることを確認してください。サポートされている計算機能を持つデバイスをコンパイラーに指定してもこのようなエラーが発生する場合、-Xsycl-target-backend に 32 ビット・トリプル nvptx-NVIDIA-cuda を渡していると考えられます。nvptx-NVIDIA-cuda トリプルは、ターゲット機能のビルトインをコンパイルできず、DPC++ では公式にサポートされていません。64 ビット・トリプル nvptx64-NVIDIA-cuda は、最近の NVIDIA* デバイスをすべてサポートするため、こちらを使用することを推奨します。

コンパイラーの警告: “CUDA version is newer than the latest supported version (CUDA* のバージョンがサポートされるバージョンよりも新しいです)”

このリリースでは、使用する CUDA* のバージョンによっては、コンパイラーが次のような警告を出力することがあります。

clang++: warning: CUDA version is newer than the latest supported version 11.5 [-Wunknown-cuda-version]

通常、この警告は無視してもかまいません。DPC++ は、最新の CUDA* でサポートされる機能を使用しないかもしれませんが、大部分のシナリオでは問題なく動作するはずです。

カーネル起動時のリソース不足

起動時に要求されたリソースが多すぎる

  • CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES:
    このエラーは、必要なリソースが不足しているため起動できなかったことを示します。

考えられる理由

  • デバイスの work-item (CUDA* スレッド) の最大数を超えています。

  • デバイスの最大 work-group サイズ (CUDA* ではスレッドブロック) を超えています。

  • カーネルのリソース (レジスターや共有メモリー) がデバイスの能力を超えています。

デバイスの能力をチェックしてこれらの可能性を検証し、制限要因を考慮してカーネル起動を構成することで問題を解決できます。

ただし、カーネル起動の最大 work-group サイズは、デバイスの潜在的な能力とは必ずしも同じではないため、カーネルが使用するレジスター数を理解する必要があります。大規模な work-group でレジスターの負荷が高くなると、ハードウェアで利用できるレジスター数を超えて、無効なカーネル起動が発生する可能性があります。詳細については、technical-specifications-per-compute-capability (英語) を参照してください。

利用可能なレジスターが不足している

CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES エラーは、CUDA* ブロックごとに使用するレジスターが多すぎることが原因で発生する可能性があります。

この問題を検出するには、コンパイル時に -Xcuda-ptxas --verbose オプションを指定することで、ptxas によってカーネルに割り当てられるレジスター数を簡単にチェックできます。これにより冗長モードが有効になり、バイナリー内のカーネルによるレジスターの使用状況を含むコード生成統計が出力されます。

ptxas の詳細出力の例:

ptxas info    : Compiling entry function 'my_kernel' for 'sm_75'
ptxas info    : Function properties for my_kernel
    8192 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 100 registers, 256 bytes cmem[0], 1512 bytes cmem[2]

最近、DPC++ ランタイムもこのケースを検出して、カーネルがハードウェアで利用可能なレジスター数を超えたことを示す詳細な例外をスローできるようになり、カーネルが実際に使用するレジスター数と失敗した起動構成の work-group に関する情報を提供します。

このタイプのエラーは、DPC++ の ERROR_INVALID_WORK_GROUP_SIZE エラーコードにマップされます。

Exceeded the number of registers available on the hardware.
    The number registers per work-group cannot exceed 65536 for this kernel on this device.
    The kernel uses 100 registers per work-item for a total of 1024 work-items per work-group.
-54 (PI_ERROR_INVALID_WORK_GROUP_SIZE)

通常、カーネルがマルチプロセッサーで利用可能なレジスター数を超えた場合、work-group サイズを小さくして、CUDA* ブロックで実行するスレッド数を減らすことが推奨されます。

ただし、これが望ましい解決策ではない場合、レジスター・プレッシャーを下げて特定のしきい値でスピルが行われるようにコンパイラーに指示することもできます。これにより、スレッドブロックのサイズを小さくすることなく起動を成功させることができます。DPC++ では次のように実現します。

  1. ターゲットデバイスの CUDA* アーキテクチャーまたは SM /Compute Capability を指定します。
    例えば、NVIDIA* GeForce* RTX 3060/TI の場合、次のように指定します。
    -Xsycl-target-backend --cuda-gpu-arch=sm_86

  2. コンパイルコマンドで -Xcuda-ptxas --maxrregcount=<N> オプションを指定して、カーネル内のレジスターを制限するよう PTX バックエンドに指示します。

注意
-Xcuda-ptxas --maxrregcount コンパイラー・オプションによってカーネルのレジスター使用を制限すると、残りのレジスターが DRAM に排出されパフォーマンスに影響する可能性があります。

プラットフォーム/アーキテクチャー間で移植されたコードの sub-group サイズの問題

カーネル属性 reqd_sub_group_size を使用して特定の sub-group サイズを設定し、その後、異なるプラットフォームに移植するか、元のアーキテクチャーとは異なるアーキテクチャーで実行されるコードについて考えてみます。このような場合、要求される sub-group サイズがプラットフォーム/アーキテクチャーでサポートされないと、実行時に次のようなエラーがスローされます。

Sub-group size x is not supported on the device

CUDA* プラットフォームでは、単一の sub-group サイズのみがサポートされるため、次の警告が出力されます。

CUDA requires sub_group size 32

そして、ランタイムは要求された sub-group サイズに代わって sub-group サイズ 32 を適用します。reqd_sub_group_size カーネル属性は、複数の sub-group サイズをサポートするプラットフォーム/アーキテクチャー向けに設計されています。一部の SYCL* コードは、異なる sub-group サイズ間では移植できないことに注意してください。例えば、sub-group 集合の reduce_over_group の結果は、sub-group サイズに依存します。異なる sub-group サイズを使用するプラットフォーム/アーキテクチャー間で移植できるコードを作成する場合、次のいずれかを考慮する必要があります。

  • 結果が sub-group サイズに依存しないよう、移植可能な方法でコードを記述します。

  • コードの sub-group サイズに依存する部分については、sub-group サイズの違いを考慮して、プラットフォーム/アーキテクチャーごとに異なるバージョンを用意します。

© Codeplay Software Ltd.

SYCL and SPIR are trademarks of the Khronos® Group. NVIDIA and CUDA are registered trademark of NVIDIA Corporation. AMD is a registered trademark of Advanced Micro Devices, Inc. Intel is a trademark of Intel Corporation in the U.S. and/or other countries. Linux is the registered trademark of Linus Torvalds in the U.S. and other countries.

関連記事