特殊化定数変数

特殊化定数は、ホストコードによって設定され、デバイスコードで使用される SYCL* プログラム内の変数です。

特殊化定数は、デバイスコード用のオンラインのジャストインタイム (JIT) コンパイラーの定数です。タイル化された行列乗算カーネルの最適なタイルサイズなどの値は、ハードウェアに依存する可能性があり、より良いコードを生成するため特殊化定数を介して表現することができます。

oneAPI では、Codeplay* からの提案 (英語) に基づき、特殊化定数の実験的な実装を提供しています。

将来のバージョンでは、この提案は SYCL* 2020 仕様 (英語) に取って代わられる可能性があります。

特殊化定数は、C++ の型名で識別されます。カーネルと同様に、値は program クラス API (set_spec_constant) で設定され、プログラムがビルドされると固定となります。次の例は、同じカーネル内で異なる値の特殊化定数を使用する方法を示します。

for (int i = 0; i < n_sc_sets; i++) {
  cl::sycl::program program(q.get_context());
  const int *sc_set = &sc_vals[i][0];
  cl::sycl::ONEAPI::experimental::spec_constant<int32_t, SC0> sc0 =
      program.set_spec_constant<SC0>(sc_set[0]);
  cl::sycl::ONEAPI::experimental::spec_constant<int32_t, SC1> sc1 =
      program.set_spec_constant<SC1>(sc_set[1]);

  program.build_with_kernel_type<KernelAAA>();

  try {
    cl::sycl::buffer<int, 1> buf(vec.data(), vec.size());

    q.submit([&](cl::sycl::handler &cgh) {
      auto acc = buf.get_access<cl::sycl::access::mode::write>(cgh);
      cgh.single_task<KernelAAA>(
          program.get_kernel<KernelAAA>(),
          [=]() {
            acc[i] = sc0.get() + sc1.get();
          });
    });
  } catch (cl::sycl::exception &e) {
    std::cout << "*** Exception caught: " << e.what() << "\n";
    return 1;
  }
  ...
}

上記の例では、特殊化定数 SC0SC1 の値はループの反復ごとに変わります。以降のコードで program クラスのインスタンスを再度作成して新しい値を設定し、program::build_with_kernel_type でリビルドする必要があります。JIT コンパイラーは、デバイスコード内の sc0.get()sc1.get() を、対応する定数値 (sc_vals[i][0]sc_vals[i][1]) に置き換えます。

Ahead-Of-Time (AOT) コンパイルでコンパイルされたプログラムで特殊化定数を使用することもできます。その場合、特殊化定数は (C++ 標準 (英語) で定義されているように) その型のデフォルト値になります。

制限