•SYCLomatic オープンソース・プロジェクト
https://www.intel.com/content/www/us/en/developer/articles/technical/syclomatic-new-cuda-to-sycl-code-migration-tool.html#inpage-nav-undefined
•SYCLomatic ツールの仕組み
https://www.intel.com/content/www/us/en/developer/articles/technical/syclomatic-new-cuda-to-sycl-code-migration-tool.html#inpage-nav-1
•例: CUDA のベクトル加算を SYCL に変換する
https://www.intel.com/content/www/us/en/developer/articles/technical/syclomatic-new-cuda-to-sycl-code-migration-tool.html#inpage-nav-3
【コード記述に関する最新情報を入手】
登録
https://software.seek.intel.com/developer-products-sign-up
Kent Moffat
インテル コーポレーション、製品担当シニア・マネージャー
CPU、GPU、FPGA など複数のアーキテクチャーにわたって開発者の生産性をパフォーマンスと効率の両面で向上するには、目の前のタスクに最適なハードウェアを選択できる、統一されたプログラミング・モデルが不可欠です。開発者は、標準規格に準拠した上で拡張性もある、オープン・スタンダードで高水準のヘテロジニアス・プログラミング言語を必要としています。生産性向上と同時に、異なるアーキテクチャーにわたる一貫したパフォーマンスも必要です。SYCL は、C++ をベースとした Khronos Group の規格で、C++ の機能を拡張してマルチアーキテクチャーと不連続のメモリー構成に対応し、上記の課題に対処します。
できるだけ手間をかけずに SYCL を導入して、ゼロから SYCL 開発を開始しなくてすむように、開発者の多くは既存の CUDA GPU コードを移植したいと考えるのではないでしょうか。インテルはこれまでに、インテル® DPC++ 互換性ツールを使用した CUDA から SYCL への移植について記事を公開しています。このツールはインテル® oneAPI ベース・ツールキットに含まれており、インテルのテクニカル・コンサルティング・エンジニアによるサポートの対象です。
＜SYCLomatic オープンソース・プロジェクト＞
この互換性ツールは現在、開発者の要望に応え、「SYCLomatic」と呼ばれるオープンソース・プロジェクトとしてリリースされています。数多くの組織がこのツールを有効に活用し、中には機能拡張やカスタマイズを行ってニーズに合わせてチューニングしたいと考える企業 / 団体もありました。その 1 つがアルゴンヌ国立研究所です。
「CRK-HACC は、精力的に開発が進められている宇宙論的 N 体シミュレーション・コードです。Aurora への準備として、20 以上の [CUDA] カーネルを SYCL に移植しましたが、インテル® DPC++ 互換性ツールによって短時間で完了できました。現行バージョンではファンクターへの移植がサポートされていないので、シンプルな Clang ツールを記述し、出力された SYCL ソースコードをニーズに合うようにリファクタリングしています。オープンソースの SYCLomatic プロジェクトにより、これまでの成果物を統合して、さらに強力なソリューションを実現するとともに、移植オプションの一部としてファンクターを利用できるようにしていく予定です」 ― アルゴンヌ国立研究所 (https://www.anl.gov/)、宇宙論物理学 & 先進コンピューティング部門、ハードウェア / ハイブリッド加速宇宙論コード (HACC) 担当 Steve (Esteban) Rangel 氏
SYCLomatic プロジェクトは、LLVM 例外付き Apache 2.0 ライセンスのもと GitHub でホスティングされている、開発者の貢献やフィードバックによって CPU、GPU、FPGA と異なるアーキテクチャーを横断したオープンなヘテロジニアス開発を推進するコミュニティーとなっています。GitHub ポータルには、このプロジェクトに技術的貢献をする場合の手順を示したガイド「contributing.md」もあります。このツールをさらに進化させるために、開発者からのフィードバックや貢献が奨励されます。このオープンソースのプロジェクトによって、コミュニティーでのコラボレーションが可能となり、SYCL 規格の採用が促進されます。これは単一ベンダー独自のエコシステムから開発者を解放する重要なステップです。SYCLomatic への改良点は、インテル® DPC++ 互換性ツール製品にも組み込まれます。
＜SYCLomatic ツールの仕組み＞
SYCLomatic は CUDA コードを SYCL に移植する開発者を支援し、多くの場合、CUDA コードの 90 ～ 95% が自動的に SYCL コードに移植されます。1 このプロセスを完了するには、開発者が残りのコーディングを手動で行ってから、ターゲット・アーキテクチャーで必要なパフォーマンス・レベルにチューニングします (図 1)。
図 1. SYCLomatic のワークフロー
＜コード移植の成功事例＞
多数の研究機関とインテルのお客様が、SYCLomatic と同じテクノロジーを備えたインテル® DPC++ 互換性ツールを使用して、複数ベンダーのアーキテクチャー上で CUDA コードから SYCL (または oneAPI の SYCL 実装である Data Parallel C++) への移植に成功しています。例えば、ストックホルム大学の GROMACS 2022、ベルリン・ツーゼ情報技術研究所 (ZIB) の easyWave、Samsung Medison、Bittware などです (ほかの例については、oneAPI DevSummit のコンテンツを参照してください)。また、アルゴンヌ国立研究所の Aurora スーパーコンピューター、ライプニッツ・スーパーコンピューティング・センター (LRZ)、GE Healthcare など、複数のお客様が現行および発売間近のインテル® Iris® Xe アーキテクチャー搭載 GPU 上でコードをテストしています。
https://www.intel.com/content/www/us/en/newsroom/news/gromacs-oneapi-aid-open-source-drug-discovery.html
http://www.ixpug.org/resources/download/from-cuda-to-dpc-back-to-nvidia-gpus-and-fpgas-an-oneapi-case-study-with-the-tsunami-simulation-easywave
http://www.youtube.com/watch?v=XBJVr5MzfBM
http://www.youtube.com/watch?v=8dNrStoJMwE
http://www.oneapi.io/events/
http://www.alcf.anl.gov/news/intel-s-oneapi-provides-tools-prepare-code-aurora
http://www.lrz.de/presse/ereignisse/2021-05-04-SuperMUC-NG-Phase-2_ENG/
http://www.youtube.com/watch?v=UJ2SclsFiIA
＜例: CUDA のベクトル加算を SYCL に移植する＞
実際の移植プロセスを説明するため、ここでは CUDA で実装されたシンプルなベクトル加算を使用し、SYCLomatic が生成するコードを詳細に見ていきます。主に注目するのは、CUDA と SYCL の違いが最も大きいコード・セクションです。このタスクには SYCLomatic とインテル® oneAPI ベース・ツールキットのインテル® oneAPI DPC++ / C++ コンパイラーを使用します。ツールキットをインストールするには、インテル® oneAPI インストール・ガイドの手順に従ってください。次のワークフローに従って、既存の CUDA アプリケーションを SYCL に移植します。
https://www.intel.com/content/www/us/en/developer/articles/guide/installation-guide-for-oneapi-toolkits.html
1.intercept-build ユーティリティーを使用して、Makefile によって発行されたコマンドをインターセプトし、JSON 形式のコンパイル・データベース・ファイルに保存します。単一ソースのプロジェクトの場合、この手順はオプションです。
2.SYCLomatic を使用して CUDA コードを SYCL に移植します。
3.生成されたコードが正しいか確認し、警告メッセージが表示された場合は、手動で移植を完了します。警告の修正方法については、インテル® DPC++ 互換性ツール開発ガイド & リファレンスを確認してください。
https://www.intel.com/content/www/us/en/develop/documentation/intel-dpcpp-compatibility-tool-user-guide/top/diagnostics-reference.html
4.インテル® oneAPI DPC++ / C++ コンパイラーを使用してコードをコンパイルし、プログラムを実行して、出力をチェックします。
その後、インテル® VTune™ プロファイラーなど、インテル® oneAPI の分析 / デバッグのツールを使用して、さらにコードを最適化できます。
例としてベクトル加算を取り上げてみましょう。ベクトル加算では、ベクトル A と B の要素を加算してベクトル C に格納します。CUDA カーネルは、これを次のように計算します。
CUDA ではスレッドの集まりがスレッドブロックです。これは SYCL のワークグループに相当しますが、スレッドのインデックス計算が異なります。CUDA では組込み変数を使用してスレッドを識別します (上記のコードで idx 変数をどのように計算しているか見てください)。SYCL に移植すると、同じカーネルが次のようになります。
CUDA のスレッドと同様に、SYCL のワークアイテムはグローバル空間でグローバル識別子、ワークグループ内でローカル識別子を持ちます。これらの識別子は nd_item 変数から取得可能です。そのため、グローバル識別子を明示的に計算する必要はありません。しかし、この例では SYCL でその計算を行う方法を示しているので、CUDA の組込み変数と類似していることが分かります。ここで nd_items に注目してください。CUDA で dim3 型を使用しているので、3 次元になっています。このコンテキストでは、nd_items を 1 次元にすることができます。それには、ワークアイテムをベクトルの各要素にマッピングします。CUDA カーネルの実行に必要なのはブロックサイズとブロック数の設定ですが、SYCL では実行範囲を指定しなければなりません。そのため、下記のコードのように、グローバル範囲とローカル範囲を組み合わせた nd_range 変数を使用します。グローバル範囲はワークアイテムの総数、ローカル範囲はワークグループのサイズです。
SYCL カーネルを呼び出すには、parallel_for と実行範囲を使用してカーネルをキューに渡します。各ワークアイテムがカーネルを呼び出すのは 1 度です。このコンテキストでは、各ベクトル要素に対して同数のワークアイテムがあります。コードは次のようになります。
これまでカーネルを実装し、実行する方法を説明してきましたが、カーネルを実行する前に、メモリー割り当てを考え、データをデバイスにコピーする必要があります。
1.まず、ホスト側でオペランドベクトル用のメモリーを割り当て、初期化します。
2.次に、同じことをデバイス上でも行います。CUDA では cudaMalloc ルーチンを使用します。DPCT はこのルーチンをデフォルトで malloc_device へ移植し、これには統合共有メモリー (USM) を使用します。
3.次に、memcpy コマンドを使用してベクトルをホストのメモリーからデバイスにコピーします。
これらの手順が完了したら、カーネルを実行します。実行後、結果をホストにコピーし、結果が正しいかチェックします。最後に、ホストとデバイスでそれぞれ free、sycl::free を呼び出し、メモリーを解放します。
＜まとめ＞
Khronos SYCL C++ 規格は、複数のアーキテクチャーで動作するヘテロジニアス・コードを開発するためのオープンな手段です。新しいオープンソース・プロジェクトである SYCLomatic は、CUDA から SYCL へのコード移植ツールを提供し、すでにリリースされているインテル® DPC++ 互換性ツールと同様のメリットがあります。現在、このツールの改良やニーズに合わせたチューニングに、だれでも貢献することができます。ぜひ挑戦してみてください。
＜開発者向けリソース＞
•Migrate from CUDA to SYCL
https://www.intel.com/content/www/us/en/developer/tools/oneapi/training/migrate-from-cuda-to-cpp-with-sycl.html
•CUDA から SYCL へのアプリケーション・カタログ
https://www.intel.com/content/www/us/en/developer/tools/oneapi/training/migrate-cuda-to-sycl-library.html
•Data Parallel C++: Programming Accelerated Systems Using C++ and SYCL
https://link.springer.com/book/10.1007/978-1-4842-9691-2
•GitHub の SYCLomatic プロジェクト | Contributing.md ガイド
http://github.com/oneapi-src/SYCLomatic/blob/SYCLomatic/CONTRIBUTING.md
•Comparing Programming models: CUDA and SYCL
http://www.codeproject.com/Articles/5324827/Comparing-Programming-models-SYCL-and-CUDA
•SYCL の詳細
http://www.khronos.org/sycl/
•oneAPI の仕様
http://www.oneapi.io/
•インテル® oneAPI ツールキット
https://www.intel.com/content/www/us/en/developer/tools/oneapi/toolkits.html
•インテル® デベロッパー・クラウド: インテルに最適化されたソフトウェアを使用して、最新のインテル® Xeon® プロセッサーや GPU コンピューティング上で AI 開発を加速できます。
https://www.intel.com/content/www/us/en/developer/tools/devcloud/services.html
【関連情報】
＜記事＞
•Migrate CUDA to DPC++ Using Intel® DPC++ Compatibility Tool
https://www.intel.com/content/www/us/en/developer/articles/technical/migrating-cuda-to-dpc-using-dpc-compatibility-tool.html
•CUDA, SYCL, Codeplay, and oneAPI: A Functional Test Walkthrough
https://www.intel.com/content/www/us/en/developer/articles/technical/cuda-sycl-oneapi-functional-test-walkthrough.html
•Migrating the Jacobi Iterative Method from CUDA to SYCL
https://www.intel.com/content/www/us/en/developer/articles/technical/cuda-sycl-migration-jacobi-iterative-method.html
•Free Your Software from Vendor Lock-in using SYCL and oneAPI
https://www.intel.com/content/www/us/en/developer/articles/technical/free-software-from-vendor-lock-in-with-sycl-oneapi.html
＜オンデマンドのウェビナー＞
•Migrate Your Existing CUDA Code to Data Parallel C++
https://www.intel.com/content/www/us/en/developer/videos/migrate-existing-cuda-code-to-data-parallel-c.html
•Optimize Edge Compute Performance by Migrating CUDA to DPC++
https://www.intel.com/content/www/us/en/developer/videos/optimize-edge-compute-performance-migrate-cuda-dpc.html
•SYCL* Essentials: Introduction to oneAPI Heterogeneous Computing
https://www.intel.com/content/www/us/en/developer/videos/sycl-essentials-introduction-workshop.html
＜製品とパフォーマンスに関する情報＞
2021年9月時点でのインテルによる推定値。Rodinia、SHOC、PENNANT など、70 の HPC ベンチマークとサンプルのセットを使用して測定した結果に基づきます。結果は異なる場合があります。
2 性能は、使用状況、構成、その他の要因によって異なります。詳細については、https://www.Intel.com/PerformanceIndex/ (英語) を参照してください。