DPC++ への Codeplay の貢献により NVIDIA* GPU の SYCL* サポートを提供

この記事は 2020 年 2 月 3 日に Codeplay のウェブサイトで公開された「Codeplay contribution to DPC++ brings SYCL support for NVIDIA GPUs」を Codeplay の許可を得て日本語訳したものです。

Codeplay は当初から SYCL* コミュニティーの一員であり、私たちのチームは過去 5 年間インテルや Xilinx を含む半導体ベンダーと協力して SYCL* 標準を定義してきました。その間、SYCL* が開発者に提供するインターフェイスは進化し、今年は最新バージョンの標準化に関する話題を提供できると期待しています。SYCL* の実装である ComputeCpp は SYCL* 1.2.1 に準拠した最初の実装であり、私たちは非常に誇りを持っており、パフォーマンスを向上させ、サポートするデバイス数を増やすため開発を続けています。

昨年、インテルが SYCL* の開発に多大な努力を払い、CPU、GPU、FPGA 向けの DPC++ (拡張機能を備えた SYCL* 実装) を含む oneAPI (英語) を開発していることを知り興奮を覚えました。さらに Xilinx、Renesas、および Imagination Technologies による SYCL* のサポートにより、ソフトウェア開発者は SYCL* を使用してさまざまなデバイスをターゲットにできるようになりました。

ComputeCpp は OpenCL* と NVIDIA の PTX を使用して NVIDIA GPU の実験的なサポートを提供していますが、DPC++ (インテルの SYCL* 実装) は OpenCL* を使用せずに LLVM コンパイラーで NVIDIA GPU を完全にサポートする可能性を提供しています。Codeplay は、SYCL* 開発者が NVIDIA GPU をターゲットにできるようにするため、我々の実装の初期の実験的なステージをオープンソース化できることを喜ばしく思います。

この実装のコードベースは、DPC++ LLVM コンパイラー・プロジェクト (英語) で公開されています。

2021 年 3 月のアップデート: このブログの投稿時点では実装は初期段階でしたが、現在では必要とされるほとんどの機能が含まれており、LBNL と ANL とのパートナーシップ (英語) による進行中のプロジェクトで、開発作業を継続しています。

NVIDIA GPU サポートの実装方法

ComputeCpp で NVIDIA プラットフォームをサポートするために採用したアプローチは、NVIDIA OpenCL* 1.2 実装に依存し、デバイス・コンパイラーを使用して SPIR-V* の代わりに PTX (NVIDIA の中間表現) を生成することでした。このアプローチは、オープン標準を使用してさまざまなプラットフォームで実行される SYCL* を実証するには十分でしたが、私たちは NVIDIA OpenCL* 1.2 によりサポートされる機能の制限にすぐに気付きました。

DPC++ コンパイラーに追加する NVIDIA プラットフォームのサポートは、OpenCL* ではなく NVIDIA CUDA* に直接関係しています。DPC++ は、プラグイン・インターフェイス (PI) を使用して各種バックエンドをターゲットにします。インテルは、DPC++ を OpenCL* プラットフォームで実行できるようにする OpenCL* 2.2 プラグインを提供しており、Codeplay では実行時に選択可能なプラグインを実装しました (環境変数 SYCL_BE=PI_CUDA を設定)。
ネイティブ CUDA* を使用することで、DPC++ は NVIDIA の OpenCL* サポートに依存せず、より多くの機能が利用できるようになり、全体のパフォーマンスが向上する可能性があります。

さらに libclc ライブラリーにいくつかの変更を実装し、SPIR-V* から作成された PTX ビルトインをサポートしました。これにより、NVIDIA GPU で SYCL* ビルトインを使用できるようになります。

SYCL* for CUDA* アプリケーションは、ネイティブ CUDA* アプリケーションであるため、CUDA* エコシステムの既存のツールとライブラリーはすべて、このバックエンドを使用してビルドされた SYCL* アプリケーションで動作します。

プロジェクトの用途

このプロジェクトでは、システムの OpenCL* レイヤーを経由せずに SYCL* コードが NVIDIA GPU をターゲットにできます。NVIDIA GPU を使用している場合、DPC++ を起動して SYCL* アプリケーションをコンパイルできます。既存の CUDA* アプリケーションがある場合、CUDA* サポートを使用して SYCL* に段階的に移行し、最終的に CUDA* がないプラットフォームで実行できるようになります。これにより、開発者はアプリケーションを他のプラットフォームに素早く移行できます。

プロジェクトの使い方

プロジェクトの README ファイルで DPC++ 用の NVIDIA バックエンドの使い方が説明されています。これは容易に利用できますが、コンパイル時にいくつかのオプションを追加する必要があります。また、ランタイムがターゲットデバイスを確実に認識できるようデバイスセレクターを設定するコードの追加も必要です。「NVIDIA CUDA* をサポートする SYCL* ツールチェーンのビルド」 (英語) と「CUDA* プラットフォームで DPC++ ツールチェーンを使用する」 (英語) で手順が説明されています。

SYCL* では、デバイスセレクターを使用してコードを実行するターゲットデバイスを指定します。NVIDIA デバイスの場合、次のように SYCL* デバイスセレクターを指定する必要があります。

class CUDASelector : public cl::sycl::device_selector {
   public:
      int operator()(const cl::sycl::device &amp;Device) const override {
         using namespace cl::sycl::info;

         const std::string DeviceName = Device.get_info();
         const std::string DeviceVendor = Device.get_info();

         if (Device.is_gpu() &amp;&amp; (DeviceName.find(&quot;NVIDIA&quot;) != std::string::npos)) {
            return 1;
         };
         return -1;
      }
};

できることとできないこと

このプロジェクトは、Titan RTX* GPU (compute capabilities 7.5) で CUDA* 10.1 を使用して、Ubuntu* 18.04 でテストされています。SM 5.0 以降と互換性のある NVIDIA GPU を搭載した他の Linux* バージョンでも動作すると思われます。

現時点では、コンパイルされた SYCL* アプリケーションは、CUDA* または OpenCL* のどちらかをターゲットにできますが、同時に両方をターゲットにすることはできません。CUDA* バックエンド向けの SYCL* アプリケーションをビルドするには、次に示すように -fsycl-targets=nvptx64-nvidia-cuda-sycldevice オプションを指定する必要があります。

$ clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice  sycl-app.cpp -o sycl-app.exe

そして、アプリケーションを実行する際に、CUDA* SYCL* バックエンドを選択する必要があります。

$ SYCL_BE=PI_CUDA sycl-app.exe

これは、PI インターフェイスが複数のバックエンドを同時にサポートできるようになると変更されます。

また、NVIDIA OpenCL* ヘッダーが DPC++ の OpenCL* ヘッダーと競合する状況も認識されています。NVIDIA OpenCL* ヘッダーは、DPC++ ではサポートされない古いバージョンの OpenCL* 向けです。DPC++ の CMake 設定は独自のヘッダーをインストールしますが、一部の設定では NVIDIA OpenCL* ヘッダーが選択されることが判明しています。

提供される機能は NVIDIA OpenCL* 実装で現在実行できる機能と一致していますが、ベースとなるハードウェアのさらに多くの機能を利用できるよう SYCL* の拡張に取り組んでいます。

期待できるパフォーマンス

プロジェクトの初期リリースは十分に最適化されていないため、すべての状況でパフォーマンスが最適であるとは限りません。

現在、コード生成は特定の最適化パスを持っていないため、プラグイン・インターフェイスの SYCL* 仕様の CUDA* 実装は高度な CUDA* 機能を使用できません。

サポート

NVIDIA GPU のサポートに関連する問題が生じた場合、Intel/llvm リポジトリー (英語) に問題を報告してください。

今後の計画

現在の目標は、開発者が NVIDIA GPU プラットフォームで各種 SYCL* アプリケーションを実行できるように、可能な限り多くの機能を実装することです。十分に機能がカバーされていると確信できたら、コミュニティーと協力して、対処すべきパフォーマンスのボトルネックを特定していきます。

私たちの目標は、SYCL* 開発者のエクスペリエンスを向上させ、CUDA* プラットフォーム上のデバイスに対する SYCL* サポートを強化することです。Codeplay によってオープンソース化されたコンポーネントは、ComputeCpp を含むほかの実装で再利用可能であるため、NVIDIA ハードウェアのエコシステム・サポートが 2020 年中に増加すると予測されます。

プロジェクトへの貢献

私たちは、さまざまな NVIDIA GPU でコードがどのように実行されるか確認して、パフォーマンスのボトルネックがどこにあるか、そしてそれぞれのハードウェア・モデル間でどのような違いがあるかを特定することに関心があります。皆さんの結果を提供いただける場合、Intel/llvm リポジトリー (英語) から結果を送信してください。

SYCL および SPIR は、Khronos Group Inc の商標です。
Nvidia および COUDA は、NVIDIA Corporation の登録商標です。
OpenCL および OpenCL ロゴは、Apple Inc. の商標であり、Khronos の使用許諾を受けて使用しています。
Imagination Imagination Technologies は、Imagination Technologies Limited の登録商標です。
Intel、インテルは、アメリカ合衆国および / またはその他の国における Intel Corporation の商標です。
Linux* は Linus Torvalds のアメリカ合衆国および / またはその他の国における商標です。

Codeplay
CTO
Ruyman Reyes