明示的 SIMD SYCL* 拡張

oneAPI は、低レベルのインテル® GPU プログラミング向けに明示的 SIMD SYCL* 拡張 (ESIMD) を提供しています。

ESIMD は、インテルの GPU 命令セット・アーキテクチャー (ISA) に似た API を提供していますが、明示的にベクトル化されたデバイスコードを記述できます。これにより、生成されるコードをより細かく制御できるようになり、コンパイラーによる最適化への依存度を下げることができます。

仕様 (英語) および 実行可能なサンプルコード (英語) は GitHub* 上で公開されています。

この拡張は、現在開発中であり、API は変更される可能性があります。以下に示すように、現在、いくつかの制限があります。

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::INTEL::gpu;

    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 コードのコンパイルと実行

ESIMD 拡張を使用するコードをコンパイルするには、特別なコンパイラー・オプション -fsycl-explicit-simd を使用する必要があります。

$ clang++ -fsycl -fsycl-explicit-simd vadd_usm.cpp

生成される実行形式は、インテル® UHD グラフィックス 600 以降などのインテル® GPU ハードウェアでのみ実行できます。実行するには、環境変数 SYCL_BE=PI_OPENCLSYCL_PROGRAM_COMPILE_OPTIONS=-vc-codegen を使用する必要があります。次に例を示します。

$ SYCL_BE=PI_OPENCL SYCL_PROGRAM_COMPILE_OPTIONS=-vc-codegen ./a.out

制約事項

このセクションでは、ESIMD 拡張を使用する際の主な制約事項の一覧を示します。

いくつかの拡張はコンパイラーによって強制されないため、プログラムの動作が未定義になる可能性があります。

ESIMD でサポートされていない機能

サポートされていない標準 SYCL* API

  • ローカルアクセサー
  • ほとんどの画像 API
  • 特殊化定数
  • sycl::accessor::get_pointer() によって返される生ポインターを介したメモリーアクセス

その他の制限

  • インテル® GPU デバイスのみサポートされます。
  • ほとんどの場合、標準カーネルと ESIMD データ並列 C++ (DPC++) カーネルは、同じアプリケーションで共存できません。