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*.
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.