Developer Guide and Reference

Contents

Specialization Constant Variables

The specialization constant is a variable in a SYCL* program set by the host code and used in the device code.
The specialization constant is a constant for the online, Just-in-Time (JIT) compiler for the device code. Values such as optimal tile size in a tiled matrix multiplication kernel may depend on your hardware and can be expressed via a specialization constant for better code generation.
oneAPI provides experimental implementation of specialization constants based on the proposal from Codeplay*.
In future versions, this proposal may be superseded by the SYCL 2020 specification.
A specialization constant is identified by a C++ type name. Similar to a kernel, its value is set with a
program
class API (
set_spec_constant
) and is frozen once the program is built. The following example shows how different values of a specialization constant can be used within the same kernel:
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; } ... }
In the example above, the values of specialization constants
SC0
and
SC1
are changed on every loop iteration. After that in your code, you must recreate a
program
class instance, set new values, and rebuild with
program::build_with_kernel_type
. The JIT compiler replaces
sc0.get()
and
sc1.get()
within the device code with the corresponding constant values (
sc_vals[i][0]
and
sc_vals[i][1]
).
You can use specialization constants in programs compiled with Ahead-Of-Time (AOT) compilation. In that case, a specialization constant takes a default value for its type (as specified by the C++ standard.
  • The implementation does not support the template
    <unsigned NID> struct spec_constant_id
    API design for interoperability with OpenCL*. Set specialization constants in SYCL programs originating from external SPIR-V* modules that are wrapped by OpenCL program objects. In SPIR-V/OpenCL, specialization constants are identified by an integer, which is modeled by the
    spec_constant_id
    class.
  • Only primitive numeric types are supported.

Product and Performance Information

1

Performance varies by use, configuration and other factors. Learn more at www.Intel.com/PerformanceIndex.