oneAPI は、低レベルのインテル® GPU プログラミング向けに明示的 SIMD SYCL* 拡張 (ESIMD) を提供しています。
ESIMD は、インテルの GPU 命令セット・アーキテクチャー (ISA) に似た API を提供していますが、明示的にベクトル化されたデバイスコードを記述できます。この明示的な有効化により、生成されるコードをより細かく制御できるようになり、コンパイラーによる最適化への依存度を下げることができます。
仕様 (英語)、ドキュメント化された ESIMD API ヘッダー (英語)、および 実行可能なサンプルコード (英語) は GitHub* 上で公開されています。
ESIMD カーネルと関数は常にサブグループ・サイズ 1 を必要とします。つまり、コンパイラーはサブグループ内のワークアイテムにまたがるベクトル化を提供しません。代りに、コード内で明示的にベクトル化を表現する必要があります。次の例は、2 つの配列の要素を加算して結果を 3 つ目の配列に書き込みます。
float *A = static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt)); float *B = static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt)); float *C = static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt)); for (unsigned i = 0; i < Size; ++i) { A[i] = B[i] = i; } // 多くのワークアイテムが必要。各ワークアイテムは // ベクトル長 (VL) の要素のデータを処理する cl::sycl::range<1> GlobalRange{Size / VL}; // 各ワークグループのワークアイテムの数。 cl::sycl::range<1> LocalRange{GroupSize}; cl::sycl::nd_range<1> Range(GlobalRange, LocalRange); auto e = q.submit([&](handler &cgh) { cgh.parallel_for<class Test>(Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::experimental::esimd; int i = ndi.get_global_id(0); simd<float, VL> va = block_load<float, VL>(A + i * VL); simd<float, VL> vb = block_load<float, VL>(B + i * VL); simd<float, VL> vc = va + vb; block_store<float, VL>(C + i * VL, vc); }); });
上記の例では、parallel_for に渡されるラムダ関数に特別な属性 SYCL_ESIMD_KERNEL が付けられています。この属性は、カーネルが ESIMD ベースであり、カーネル内で ESIMD API を使用できることをコンパイラーに伝えます。ここでは、simd オブジェクトと block_load/block_store 組込み関数が使用されています。これらは ESIMD 拡張でのみ利用できます。
実行可能なサンプルコードは、GitHub リポジトリー (英語) にあります。
ESIMD 拡張を使用するコードのコンパイルと実行は、標準 SYCL* を使用するコードのコンパイルと実行と同じです。
$ dpcpp vadd_usm.cpp
生成される実行形式 ($./a.out) は、インテル® UHD グラフィックス 600 以降などのインテル® GPU ハードウェアでのみ実行できます。DPC++ ランタイムは ESIMD カーネルを自動的に認識して実行をディスパッチするため、追加のセットアップは不要です。
このセクションでは、ESIMD 拡張を使用する際の主な制約事項の一覧を示します。