From 77cb1570b1759e4b33d6c72d650db93d32050852 Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Wed, 4 Dec 2019 08:18:33 +0000 Subject: [PATCH 1/8] Updated spec constant proposal to use NTTP. This patch extended the module proposal to add better support for spec constant. It also introduce a different way to name spec constant values. Spec constant are named by creating an object of type `spec_id` and can be manipulated using NTTP. --- proposals/sycl_modules.md | 204 ++++++++++++++++++++++++++++++++------ 1 file changed, 176 insertions(+), 28 deletions(-) diff --git a/proposals/sycl_modules.md b/proposals/sycl_modules.md index 90c1ca4..6cbe385 100644 --- a/proposals/sycl_modules.md +++ b/proposals/sycl_modules.md @@ -11,7 +11,7 @@ | Available | N/A | | Reply-to | Ruyman Reyes | | Original author | Ruyman Reyes | -| Contributors | Aksel Alpay , Victor Lomuller , Toomas Remmelg , Morris Hafner | +| Contributors | Aksel Alpay , Victor Lomuller , Toomas Remmelg , Morris Hafner , Roland Schulz | **The contents of this proposal are to encourage early feedback and are not ratified by the Khronos Group so do not constitute a formal specification.** @@ -106,6 +106,12 @@ in the scope of the current wording for this proposal to address the ## Revisions +v 0.10: + +* Rework of the specialization constant proposal + * Add `spec_id` and `spec_constant` definitions + * Add API to manipulate a Module spec constants + v 0.9: * Added wording to clarify `device_image` relationship with `module_state` @@ -306,7 +312,8 @@ Specialization is intended for constant objects that will not have known constan until after initial generation of a module in an intermediate representation format (e.g. SPIR-V). When a module contains *specialization constants*, the method *`sycl::module::has_spec_constant()`* returns `true`. -Note that, although a module may expose *specialization constants*, not all *device images* inside the module may support them. +Note that, although a module may expose *specialization constants*, not all *device images* inside the module may natively support them, +the method *`sycl::module::native_spec_constant()`* returns `true` if the module is capable of using *specialization constants* as immediate value when compiling the module image. When *specialization constants* are available but the user doesn't set values on them, the default values are used. The default values follow C++ initialization rules whenever not specified by users. @@ -377,14 +384,38 @@ class module { */ bool has_spec_constant() const noexcept; - // Available if module_state::input - /* set_spec_constant. - * If ID is a spec constant residing in the current module, returns a - * spec_constant immutable object to represent the binding of the value - * to the given module. - */ - template - spec_constant set_spec_constant(T cst); + /** + * Returns true if the current module can support specialization constants natively. + * + */ + bool native_spec_constant() const noexcept; + + // Available if module_status::executable + // Retrive the value of the spec constant associated with this module. + // Call only valid from the host. + template& s> + spec_constant get_spec_constant(handler&); + + // Available if module_status::executable + // Only from C++17 + // Set the value of the spec constant. + // Call only valid from the host. + template + spec_constant::type, s> + get_spec_constant(handler&); + + // Available if module_status::input + // Set the value of the spec constant. + // Call only valid from the host. + template& s> + void set_spec_constant(T); + + // Available if module_status::input + // Only from C++17 + // Set the value of the spec constant. + // Call only valid from the host. + template + void set_spec_constant(typename std::remove_reference_t::type); }; } // namespace sycl @@ -707,23 +738,46 @@ value to the specialization constant. The program can then be build with `build_with_kernel_type` to create a specialized kernel. +### Specialization constant types + On the new, type-safe, module approach, specialization constants are -associated with the OpenCL backend. -Other backends (e.g. Vulkan) may offer the same functionality, so the -specialization constant type is defined as an (experimental) general +typically associated with SPIR-V module images but are not limited to such image. +The specialization constant type is defined as an (experimental) general SYCL type as below: ```cpp namespace sycl { namespace experimental { -template +template +class spec_id { +private: + // Implementation defined constructor. + spec_id(const spec_id&) = delete; + spec_id(spec_id&&) = delete; + +public: + + using type = T; + + // Argument `Args` are forwarded to an underlying T Ctor. + // This allow the user to setup a default value for the spec_id instance. + // The initialization of T must be evaluated at compile time to be valid. + template + explicit constexpr spec_id(Args&&...); +}; + +template & s> class spec_constant { private: // Implementation defined constructor. spec_constant(/* Implementation defined */); + spec_constant(spec_constant&&) = delete; + public: - spec_constant(); + using type = T; + + spec_constant(const spec_constant&) = default; T get() const; // explicit access. operator T() const; // implicit conversion. @@ -733,40 +787,134 @@ public: } // namespace sycl ``` +`spec_id` is used by the *SYCL implementation* to *reference* constant whose +concrete value will only be known during the application execution. +To create a new `spec_id` in a module, the user creates a `spec_id` object with an *automatic* or *static* storage duration in the namespace or class scope. +The user then uses references to this object to bind the `spec_constant` value to the *identifier*. + +`spec_id` objects must be forward declare-able and cannot be moved or copied. + `spec_constant` is a *placeholder type* used by the *SYCL implementation* to inject the real value of the specialization constant at runtime. -It also allows implementations that don't support specialization constant -to hide the value as a parameter, keeping the lambda-capture the same -irrespectively of the support of the feature. +For the module associated with the host device, and possibly any device with non native spec constant support, the value is carried by the `spec_constant` object it-self. +For modules natively supporting specialization constant, implementations must guaranty the values returned by the usage of the spec_constant materialized into a constant. + +### Setting the value of spec constant -The interface to set values to the specialization constant -in the module is backend specific: +To modify the value of a spec constant, the following API is available. + +Given a SYCL module, a spec constant can be set by the `sycl::set_spec_constant` function. ```cpp namespace sycl { -namespace opencl { - template - spec_constant set_spec_constant(module syclModule, T cst) const; -} // namespace opencl + +template& s> +void set_spec_constant(module syclModule, T); + +} // namespace sycl +``` +If C++17 or later is enabled, then the following overload is available +```cpp +namespace sycl { + +template +void set_spec_constant(module syclModule, + typename std::remove_reference_t::type); + +} // namespace sycl +``` + +It is valid for a user to set several times a spec constant with the same `spec_id`, +only the last version of the value will be consider when building the module. +Calling this function is not allowed during the submission of a task. + +### Getting the value of spec constant in kernel + +To retrive the value of a spec constant in a SYCL kernel context, the following API is available. + +```cpp +namespace sycl { + +template& s> +T get_spec_constant(const spec_constant&); + } // namespace sycl ``` +If C++17 or later is enabled, then the following overload is available +```cpp +namespace sycl { + +template +typename std::remove_reference_t::type +get_spec_constant(const spec_constant::type, s>&); + +} // namespace sycl +``` + +This call is only valid whitin a SYCL kernel context and the `spec_constant` value must have been +created by the handler that invoked the kernel from which the `sycl::get_spec_constant` is called. + +If the module has native support for specialization constant, +the returned value becomes a constant value available for the runtime compiler to use for optimisations. + +**Note:** the value may become a constant (known numerical value), although this may only happen at runtime. +From the SYCL compiler's perspective (C++ compiler) the value is understood to be a runtime value and +for this reason it is not possible to use such value in a `constexpr`. + +### Retrieving a spec constant placeholder + +A `spec_constant` handler can be retrieved from the command group handler by either: + - Getting the placeholder for a previously set specialization constant value; + - If no specialization constant where ever set, then the *placeholder* object holds the default value of its `spec_id`. + - Setting a specialization constant value and retrieving the placeholder + - Performance notice: this function may cause the SYCL runtime to implicitly recompile the relevant SYCL module with the given value. This may have an important effect on performances and users should check their vendor recommendations on how to efficiently uses this functionality. -For simplicity, the spec constant placeholder type can also be obtained from -the command group handler: +The command group `handler` class is extended with the following API: ```cpp namespace sycl { class handler { // ... public: - template - spec_constant get_spec_constant(); + + template& s> + spec_constant set_spec_constant(T); + + template& s> + spec_constant get_spec_constant(); +// ... +}; + +} // namespace sycl +``` +If C++17 or later is enabled, then the following overloads are available +```cpp +namespace sycl { + class handler { +// ... +public: + + // Same as template& s> spec_constant set_spec_constant(T) + template + spec_constant::type, s> + set_spec_constant(typename std::remove_reference_t::type); + + // Same as template& s> spec_constant get_spec_constant() + template + spec_constant::type, s> + get_spec_constant(); // ... }; } // namespace sycl ``` +| Member function | Description | +|-----------------------------------------------------------------------------|-------------------------------------------------------------------------------------------------------------------------------------------| +| `template& s> spec_constant set_spec_constant(T)` | Set the given value as the specialization constant value in the underlying module and returns a `spec_constant` handler usable in kernel. If the user specify to the handle a specific module to use, then the value given to `set_spec_constant` must match the the one used to build the module. In case of mismatch, the SYCL runtime raises an error. | +| `template& s> spec_constant get_spec_constant()` | Returns a `spec_constant` handler usable in kernel. | + + Note the value of the specialization constant depends on the module that is used, not on the placeholder object. From 7daa69967781f25c5b9be5ea703658618e08223e Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Wed, 11 Dec 2019 16:25:56 +0000 Subject: [PATCH 2/8] Remove the use of experimental namespace --- proposals/sycl_modules.md | 19 ++++++++++--------- 1 file changed, 10 insertions(+), 9 deletions(-) diff --git a/proposals/sycl_modules.md b/proposals/sycl_modules.md index 6cbe385..d5c86db 100644 --- a/proposals/sycl_modules.md +++ b/proposals/sycl_modules.md @@ -742,12 +742,10 @@ specialized kernel. On the new, type-safe, module approach, specialization constants are typically associated with SPIR-V module images but are not limited to such image. -The specialization constant type is defined as an (experimental) general -SYCL type as below: +The specialization constant type is defined as an general SYCL type as below: ```cpp namespace sycl { -namespace experimental { template class spec_id { @@ -756,6 +754,7 @@ private: spec_id(const spec_id&) = delete; spec_id(spec_id&&) = delete; + public: using type = T; @@ -769,21 +768,22 @@ public: template & s> class spec_constant { -private: - // Implementation defined constructor. - spec_constant(/* Implementation defined */); - spec_constant(spec_constant&&) = delete; - public: using type = T; + // Create an empty (invalid) spec constant. + // Remains invalid until initialized by a call + // to handler::set_spec_constant or + // to handler::get_spec_constant + spec_constant(); + spec_constant(const spec_constant&) = default; + spec_constant& operator=(const spec_constant&) = default; T get() const; // explicit access. operator T() const; // implicit conversion. }; -} // namespace experimental } // namespace sycl ``` @@ -799,6 +799,7 @@ inject the real value of the specialization constant at runtime. For the module associated with the host device, and possibly any device with non native spec constant support, the value is carried by the `spec_constant` object it-self. For modules natively supporting specialization constant, implementations must guaranty the values returned by the usage of the spec_constant materialized into a constant. + ### Setting the value of spec constant To modify the value of a spec constant, the following API is available. From fe00ff200346abcf6c8016c35db2b5bb805f53bd Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Wed, 11 Dec 2019 16:26:39 +0000 Subject: [PATCH 3/8] Add a small example --- proposals/sycl_modules.md | 53 +++++++++++++++++++++++++++++++++++++++ 1 file changed, 53 insertions(+) diff --git a/proposals/sycl_modules.md b/proposals/sycl_modules.md index d5c86db..05f11df 100644 --- a/proposals/sycl_modules.md +++ b/proposals/sycl_modules.md @@ -927,3 +927,56 @@ values on the specialization constant or (2) having the value of the specializat constant as a normal kernel parameter. If choosing (1), the original SPIR-V image should be part of the SYCL module so that the module can be specialized for non-known values. + +### Example + +The following code is showcasing how specialization constant can be used. +The code perform a simple convolution using a scaled convolution kernel. +The specialization constants are here used to inject, at runtime, +the coefficients of the kernel. + +```cpp +#include + +using namespace cl::sycl; + +using coeff_t = std::array, 3>; + +// Read coefficients from somewhere. +coeff_t get_coefficients(); + +// Identify the specialization constant. +spec_id coeff_id; + +void do_conv(buffer in, buffer out) { + queue myQueue; + + myQueue.submit([&](handler &cgh) { + auto in_acc = in.get_access(cgh); + auto out_acc = out.get_access(cgh); + + // Create a specialization constant placeholder. + spec_constant coeff = + cgh.set_spec_constant(get_coefficients()); + + cgh.parallel_for( + in.get_range(), [=](cl::sycl::item<2> item_id) { + float acc = 0; + for (int i = -1; i < 2; i++) { + if (item_id[0] + i < 0 || item_id[0] + i > in_acc.get_range()[0]) + continue; + for (int j = -1; j < 2; j++) { + if (item_id[1] + j < 0 || item_id[1] + j > in_acc.get_range()[1]) + continue; + // the underlying JIT can see all the values of the array returned by coeff.get(). + acc += coeff.get()[i + 1][j + 1] * + in_acc[item_id[0] + i][item_id[1] + j]; + } + } + out_acc[item_id] = acc; + }); + }); + + myQueue.wait(); +} +``` From 72d7513897d79853d3b3dd447b142522fb03f5f8 Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Thu, 12 Dec 2019 16:31:17 +0000 Subject: [PATCH 4/8] Restrict to one call to handler::set_spec_constant or handler::get_spec_constant per spec_id and per command group scope --- proposals/sycl_modules.md | 2 ++ 1 file changed, 2 insertions(+) diff --git a/proposals/sycl_modules.md b/proposals/sycl_modules.md index 05f11df..573a69d 100644 --- a/proposals/sycl_modules.md +++ b/proposals/sycl_modules.md @@ -915,6 +915,8 @@ public: | `template& s> spec_constant set_spec_constant(T)` | Set the given value as the specialization constant value in the underlying module and returns a `spec_constant` handler usable in kernel. If the user specify to the handle a specific module to use, then the value given to `set_spec_constant` must match the the one used to build the module. In case of mismatch, the SYCL runtime raises an error. | | `template& s> spec_constant get_spec_constant()` | Returns a `spec_constant` handler usable in kernel. | +Only one call to `set_spec_constant` or `get_spec_constant` is allowed per spec_id and per command group scope. +If more than one call per spec_id and per command group scope is made then an runtime exception is thrown. Note the value of the specialization constant depends on the module that is used, not on the placeholder object. From b29123a7b0f3110b3c3e6a0b781ca9b8faef2bd2 Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Tue, 7 Jan 2020 13:55:32 +0000 Subject: [PATCH 5/8] Add typo and some editorial fixes --- proposals/sycl_modules.md | 71 ++++++++++++++++++++------------------- 1 file changed, 37 insertions(+), 34 deletions(-) diff --git a/proposals/sycl_modules.md b/proposals/sycl_modules.md index 573a69d..493ec29 100644 --- a/proposals/sycl_modules.md +++ b/proposals/sycl_modules.md @@ -110,7 +110,7 @@ v 0.10: * Rework of the specialization constant proposal * Add `spec_id` and `spec_constant` definitions - * Add API to manipulate a Module spec constants + * Add API to manipulate a Module specialization constants v 0.9: @@ -315,8 +315,7 @@ When a module contains *specialization constants*, the method Note that, although a module may expose *specialization constants*, not all *device images* inside the module may natively support them, the method *`sycl::module::native_spec_constant()`* returns `true` if the module is capable of using *specialization constants* as immediate value when compiling the module image. When *specialization constants* are available but the user doesn't set values on them, the default values are used. -The default values follow C++ initialization rules whenever not specified by -users. +The default values follow C++ initialization rules whenever not specified by users. When *specialization constants* are not available in some images in a module, but the user sets a value to them, they are set as an additional kernel argument. @@ -391,28 +390,28 @@ class module { bool native_spec_constant() const noexcept; // Available if module_status::executable - // Retrive the value of the spec constant associated with this module. + // Retrieve the value of the specialization constant associated with this module. // Call only valid from the host. template& s> spec_constant get_spec_constant(handler&); // Available if module_status::executable // Only from C++17 - // Set the value of the spec constant. + // Set the value of the specialization constant. // Call only valid from the host. template spec_constant::type, s> get_spec_constant(handler&); // Available if module_status::input - // Set the value of the spec constant. + // Set the value of the specialization constant. // Call only valid from the host. template& s> void set_spec_constant(T); // Available if module_status::input // Only from C++17 - // Set the value of the spec constant. + // Set the value of the specialization constant. // Call only valid from the host. template void set_spec_constant(typename std::remove_reference_t::type); @@ -731,18 +730,18 @@ module. Proposal [CP015](https://github.com/codeplaysoftware/standards-proposals/blob/master/spec-constant/index.md) introduced usage of SPIR-V specialization constants in SYCL. -Specialization constants are associated with the SYCL 1.2.1 program class, +Specialization constants are associated with the SYCL 1.2.1 program class, and named using a C++ type like kernels. The program class gains a `set_spec_constant` method, that sets a runtime value to the specialization constant. -The program can then be build with `build_with_kernel_type` to create a +The program can then be build with `build_with_kernel_type` to create a specialized kernel. ### Specialization constant types On the new, type-safe, module approach, specialization constants are -typically associated with SPIR-V module images but are not limited to such image. -The specialization constant type is defined as an general SYCL type as below: +typically associated with SPIR-V module images but are not limited to such images. +The specialization constant type is defined as a general SYCL type as below: ```cpp namespace sycl { @@ -771,7 +770,7 @@ class spec_constant { public: using type = T; - // Create an empty (invalid) spec constant. + // Create an empty (invalid) specialization constant. // Remains invalid until initialized by a call // to handler::set_spec_constant or // to handler::get_spec_constant @@ -792,46 +791,47 @@ concrete value will only be known during the application execution. To create a new `spec_id` in a module, the user creates a `spec_id` object with an *automatic* or *static* storage duration in the namespace or class scope. The user then uses references to this object to bind the `spec_constant` value to the *identifier*. -`spec_id` objects must be forward declare-able and cannot be moved or copied. +`spec_id` objects must be forward declarable and cannot be moved or copied. `spec_constant` is a *placeholder type* used by the *SYCL implementation* to inject the real value of the specialization constant at runtime. -For the module associated with the host device, and possibly any device with non native spec constant support, the value is carried by the `spec_constant` object it-self. +For the module associated with the host device, and possibly any device with non native specialization constant support, the value is carried by the `spec_constant` object itself. For modules natively supporting specialization constant, implementations must guaranty the values returned by the usage of the spec_constant materialized into a constant. +### Setting the value of specialization constant -### Setting the value of spec constant +To modify the value of a specialization constant, the following API is available. -To modify the value of a spec constant, the following API is available. - -Given a SYCL module, a spec constant can be set by the `sycl::set_spec_constant` function. +Given a SYCL module, a specialization constant can be set by the `sycl::set_spec_constant` function. ```cpp namespace sycl { template& s> -void set_spec_constant(module syclModule, T); +void set_spec_constant(module &syclModule, T); } // namespace sycl ``` + If C++17 or later is enabled, then the following overload is available + ```cpp namespace sycl { template -void set_spec_constant(module syclModule, +void set_spec_constant(module &syclModule, typename std::remove_reference_t::type); } // namespace sycl ``` -It is valid for a user to set several times a spec constant with the same `spec_id`, +It is valid for a user to set several times a specialization constant with the same `spec_id`, only the last version of the value will be consider when building the module. Calling this function is not allowed during the submission of a task. -### Getting the value of spec constant in kernel +### Getting the value of specialization constant in kernel -To retrive the value of a spec constant in a SYCL kernel context, the following API is available. +To retrieve the value of a specialization constant in a SYCL kernel context, the following API is available. ```cpp namespace sycl { @@ -862,13 +862,14 @@ the returned value becomes a constant value available for the runtime compiler t From the SYCL compiler's perspective (C++ compiler) the value is understood to be a runtime value and for this reason it is not possible to use such value in a `constexpr`. -### Retrieving a spec constant placeholder +### Retrieving a specialization constant placeholder A `spec_constant` handler can be retrieved from the command group handler by either: - - Getting the placeholder for a previously set specialization constant value; - - If no specialization constant where ever set, then the *placeholder* object holds the default value of its `spec_id`. - - Setting a specialization constant value and retrieving the placeholder - - Performance notice: this function may cause the SYCL runtime to implicitly recompile the relevant SYCL module with the given value. This may have an important effect on performances and users should check their vendor recommendations on how to efficiently uses this functionality. + +- Getting the placeholder for a previously set specialization constant value; + - If no specialization constant where ever set, then the *placeholder* object holds the default value of its `spec_id`. +- Setting a specialization constant value and retrieving the placeholder + - Performance notice: this function may cause the SYCL runtime to implicitly recompile the relevant SYCL module with the given value. This may have an important effect on performances and users should check their vendor recommendations on how to efficiently uses this functionality. The command group `handler` class is extended with the following API: @@ -888,7 +889,9 @@ public: } // namespace sycl ``` + If C++17 or later is enabled, then the following overloads are available + ```cpp namespace sycl { class handler { @@ -915,8 +918,8 @@ public: | `template& s> spec_constant set_spec_constant(T)` | Set the given value as the specialization constant value in the underlying module and returns a `spec_constant` handler usable in kernel. If the user specify to the handle a specific module to use, then the value given to `set_spec_constant` must match the the one used to build the module. In case of mismatch, the SYCL runtime raises an error. | | `template& s> spec_constant get_spec_constant()` | Returns a `spec_constant` handler usable in kernel. | -Only one call to `set_spec_constant` or `get_spec_constant` is allowed per spec_id and per command group scope. -If more than one call per spec_id and per command group scope is made then an runtime exception is thrown. +Only one call to `set_spec_constant` or `get_spec_constant` is allowed per `spec_id` and per command group scope. +If more than one call per `spec_id` and per command group scope is made then an runtime exception is thrown. Note the value of the specialization constant depends on the module that is used, not on the placeholder object. @@ -964,11 +967,11 @@ void do_conv(buffer in, buffer out) { cgh.parallel_for( in.get_range(), [=](cl::sycl::item<2> item_id) { float acc = 0; - for (int i = -1; i < 2; i++) { - if (item_id[0] + i < 0 || item_id[0] + i > in_acc.get_range()[0]) + for (int i = -1; i <= 1; i++) { + if (item_id[0] + i < 0 || item_id[0] + i >= in_acc.get_range()[0]) continue; - for (int j = -1; j < 2; j++) { - if (item_id[1] + j < 0 || item_id[1] + j > in_acc.get_range()[1]) + for (int j = -1; j <= 1; j++) { + if (item_id[1] + j < 0 || item_id[1] + j >= in_acc.get_range()[1]) continue; // the underlying JIT can see all the values of the array returned by coeff.get(). acc += coeff.get()[i + 1][j + 1] * From ad02463fe97933c35a10d5bccab75f53e991613b Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Tue, 7 Jan 2020 14:58:26 +0000 Subject: [PATCH 6/8] Change spec_* to specialization_ --- proposals/sycl_modules.md | 122 +++++++++++++++++++------------------- 1 file changed, 61 insertions(+), 61 deletions(-) diff --git a/proposals/sycl_modules.md b/proposals/sycl_modules.md index 493ec29..7c1c8a6 100644 --- a/proposals/sycl_modules.md +++ b/proposals/sycl_modules.md @@ -109,7 +109,7 @@ in the scope of the current wording for this proposal to address the v 0.10: * Rework of the specialization constant proposal - * Add `spec_id` and `spec_constant` definitions + * Add `specialization _id` and `specialization _constant` definitions * Add API to manipulate a Module specialization constants v 0.9: @@ -311,9 +311,9 @@ A *SYCL module* objects may expose *specialization constants*. Specialization is intended for constant objects that will not have known constant values until after initial generation of a module in an intermediate representation format (e.g. SPIR-V). When a module contains *specialization constants*, the method -*`sycl::module::has_spec_constant()`* returns `true`. +*`sycl::module::has_specialization_constant()`* returns `true`. Note that, although a module may expose *specialization constants*, not all *device images* inside the module may natively support them, -the method *`sycl::module::native_spec_constant()`* returns `true` if the module is capable of using *specialization constants* as immediate value when compiling the module image. +the method *`sycl::module::native_specialization_constant()`* returns `true` if the module is capable of using *specialization constants* as immediate value when compiling the module image. When *specialization constants* are available but the user doesn't set values on them, the default values are used. The default values follow C++ initialization rules whenever not specified by users. When *specialization constants* are not available in some images in a module, @@ -381,40 +381,40 @@ class module { * Returns true if the current module uses specialization constants. * */ - bool has_spec_constant() const noexcept; + bool has_specialization_constant() const noexcept; /** * Returns true if the current module can support specialization constants natively. * */ - bool native_spec_constant() const noexcept; + bool native_specialization_constant() const noexcept; // Available if module_status::executable // Retrieve the value of the specialization constant associated with this module. // Call only valid from the host. - template& s> - spec_constant get_spec_constant(handler&); + template& s> + specialization_constant get_specialization_constant(handler&); // Available if module_status::executable // Only from C++17 // Set the value of the specialization constant. // Call only valid from the host. template - spec_constant::type, s> - get_spec_constant(handler&); + specialization_constant::type, s> + get_specialization_constant(handler&); // Available if module_status::input // Set the value of the specialization constant. // Call only valid from the host. - template& s> - void set_spec_constant(T); + template& s> + void set_specialization_constant(T); // Available if module_status::input // Only from C++17 // Set the value of the specialization constant. // Call only valid from the host. template - void set_spec_constant(typename std::remove_reference_t::type); + void set_specialization_constant(typename std::remove_reference_t::type); }; } // namespace sycl @@ -732,7 +732,7 @@ Proposal [CP015](https://github.com/codeplaysoftware/standards-proposals/blob/ma introduced usage of SPIR-V specialization constants in SYCL. Specialization constants are associated with the SYCL 1.2.1 program class, and named using a C++ type like kernels. -The program class gains a `set_spec_constant` method, that sets a runtime +The program class gains a `set_specialization_constant` method, that sets a runtime value to the specialization constant. The program can then be build with `build_with_kernel_type` to create a specialized kernel. @@ -747,11 +747,11 @@ The specialization constant type is defined as a general SYCL type as below: namespace sycl { template -class spec_id { +class specialization_id { private: // Implementation defined constructor. - spec_id(const spec_id&) = delete; - spec_id(spec_id&&) = delete; + specialization_id(const specialization_id&) = delete; + specialization_id(specialization_id&&) = delete; public: @@ -759,25 +759,25 @@ public: using type = T; // Argument `Args` are forwarded to an underlying T Ctor. - // This allow the user to setup a default value for the spec_id instance. + // This allow the user to setup a default value for the specialization_id instance. // The initialization of T must be evaluated at compile time to be valid. template - explicit constexpr spec_id(Args&&...); + explicit constexpr specialization_id(Args&&...); }; -template & s> -class spec_constant { +template & s> +class specialization_constant { public: using type = T; // Create an empty (invalid) specialization constant. // Remains invalid until initialized by a call - // to handler::set_spec_constant or - // to handler::get_spec_constant - spec_constant(); + // to handler::set_specialization_constant or + // to handler::get_specialization_constant + specialization_constant(); - spec_constant(const spec_constant&) = default; - spec_constant& operator=(const spec_constant&) = default; + specialization_constant(const specialization_constant&) = default; + specialization_constant& operator=(const specialization_constant&) = default; T get() const; // explicit access. operator T() const; // implicit conversion. @@ -786,29 +786,29 @@ public: } // namespace sycl ``` -`spec_id` is used by the *SYCL implementation* to *reference* constant whose +`specialization_id` is used by the *SYCL implementation* to *reference* constant whose concrete value will only be known during the application execution. -To create a new `spec_id` in a module, the user creates a `spec_id` object with an *automatic* or *static* storage duration in the namespace or class scope. -The user then uses references to this object to bind the `spec_constant` value to the *identifier*. +To create a new `specialization_id` in a module, the user creates a `specialization_id` object with an *automatic* or *static* storage duration in the namespace or class scope. +The user then uses references to this object to bind the `specialization_constant` value to the *identifier*. -`spec_id` objects must be forward declarable and cannot be moved or copied. +`specialization_id` objects must be forward declarable and cannot be moved or copied. -`spec_constant` is a *placeholder type* used by the *SYCL implementation* to +`specialization_constant` is a *placeholder type* used by the *SYCL implementation* to inject the real value of the specialization constant at runtime. -For the module associated with the host device, and possibly any device with non native specialization constant support, the value is carried by the `spec_constant` object itself. -For modules natively supporting specialization constant, implementations must guaranty the values returned by the usage of the spec_constant materialized into a constant. +For the module associated with the host device, and possibly any device with non native specialization constant support, the value is carried by the `specialization_constant` object itself. +For modules natively supporting specialization constant, implementations must guaranty the values returned by the usage of the `specialization_constant` materialized into a constant. ### Setting the value of specialization constant To modify the value of a specialization constant, the following API is available. -Given a SYCL module, a specialization constant can be set by the `sycl::set_spec_constant` function. +Given a SYCL module, a specialization constant can be set by the `sycl::set_specialization_constant` function. ```cpp namespace sycl { -template& s> -void set_spec_constant(module &syclModule, T); +template& s> +void set_specialization_constant(module &syclModule, T); } // namespace sycl ``` @@ -819,13 +819,13 @@ If C++17 or later is enabled, then the following overload is available namespace sycl { template -void set_spec_constant(module &syclModule, +void set_specialization_constant(module &syclModule, typename std::remove_reference_t::type); } // namespace sycl ``` -It is valid for a user to set several times a specialization constant with the same `spec_id`, +It is valid for a user to set several times a specialization constant with the same `specialization_id`, only the last version of the value will be consider when building the module. Calling this function is not allowed during the submission of a task. @@ -836,8 +836,8 @@ To retrieve the value of a specialization constant in a SYCL kernel context, the ```cpp namespace sycl { -template& s> -T get_spec_constant(const spec_constant&); +template& s> +T get_specialization_constant(const specialization_constant&); } // namespace sycl ``` @@ -847,13 +847,13 @@ namespace sycl { template typename std::remove_reference_t::type -get_spec_constant(const spec_constant::type, s>&); +get_specialization_constant(const specialization_constant::type, s>&); } // namespace sycl ``` -This call is only valid whitin a SYCL kernel context and the `spec_constant` value must have been -created by the handler that invoked the kernel from which the `sycl::get_spec_constant` is called. +This call is only valid whitin a SYCL kernel context and the `specialization_constant` value must have been +created by the handler that invoked the kernel from which the `sycl::get_specialization_constant` is called. If the module has native support for specialization constant, the returned value becomes a constant value available for the runtime compiler to use for optimisations. @@ -864,10 +864,10 @@ for this reason it is not possible to use such value in a `constexpr`. ### Retrieving a specialization constant placeholder -A `spec_constant` handler can be retrieved from the command group handler by either: +A `specialization_constant` handler can be retrieved from the command group handler by either: - Getting the placeholder for a previously set specialization constant value; - - If no specialization constant where ever set, then the *placeholder* object holds the default value of its `spec_id`. + - If no specialization constant where ever set, then the *placeholder* object holds the default value of its `specialization_id`. - Setting a specialization constant value and retrieving the placeholder - Performance notice: this function may cause the SYCL runtime to implicitly recompile the relevant SYCL module with the given value. This may have an important effect on performances and users should check their vendor recommendations on how to efficiently uses this functionality. @@ -879,11 +879,11 @@ namespace sycl { // ... public: - template& s> - spec_constant set_spec_constant(T); + template& s> + specialization_constant set_specialization_constant(T); - template& s> - spec_constant get_spec_constant(); + template& s> + specialization_constant get_specialization_constant(); // ... }; @@ -898,15 +898,15 @@ namespace sycl { // ... public: - // Same as template& s> spec_constant set_spec_constant(T) + // Same as template& s> specialization_constant set_specialization_constant(T) template - spec_constant::type, s> - set_spec_constant(typename std::remove_reference_t::type); + specialization_constant::type, s> + set_specialization_constant(typename std::remove_reference_t::type); - // Same as template& s> spec_constant get_spec_constant() + // Same as template& s> specialization_constant get_specialization_constant() template - spec_constant::type, s> - get_spec_constant(); + specialization_constant::type, s> + get_specialization_constant(); // ... }; @@ -915,11 +915,11 @@ public: | Member function | Description | |-----------------------------------------------------------------------------|-------------------------------------------------------------------------------------------------------------------------------------------| -| `template& s> spec_constant set_spec_constant(T)` | Set the given value as the specialization constant value in the underlying module and returns a `spec_constant` handler usable in kernel. If the user specify to the handle a specific module to use, then the value given to `set_spec_constant` must match the the one used to build the module. In case of mismatch, the SYCL runtime raises an error. | -| `template& s> spec_constant get_spec_constant()` | Returns a `spec_constant` handler usable in kernel. | +| `template& s> specialization_constant set_specialization_constant(T)` | Set the given value as the specialization constant value in the underlying module and returns a `specialization_constant` handler usable in kernel. If the user specify to the handle a specific module to use, then the value given to `set_specialization_constant` must match the the one used to build the module. In case of mismatch, the SYCL runtime raises an error. | +| `template& s> specialization_constant get_specialization_constant()` | Returns a `specialization_constant` handler usable in kernel. | -Only one call to `set_spec_constant` or `get_spec_constant` is allowed per `spec_id` and per command group scope. -If more than one call per `spec_id` and per command group scope is made then an runtime exception is thrown. +Only one call to `set_specialization_constant` or `get_specialization_constant` is allowed per `specialization_id` and per command group scope. +If more than one call per `specialization_id` and per command group scope is made then an runtime exception is thrown. Note the value of the specialization constant depends on the module that is used, not on the placeholder object. @@ -951,7 +951,7 @@ using coeff_t = std::array, 3>; coeff_t get_coefficients(); // Identify the specialization constant. -spec_id coeff_id; +specialization_id coeff_id; void do_conv(buffer in, buffer out) { queue myQueue; @@ -961,8 +961,8 @@ void do_conv(buffer in, buffer out) { auto out_acc = out.get_access(cgh); // Create a specialization constant placeholder. - spec_constant coeff = - cgh.set_spec_constant(get_coefficients()); + specialization_constant coeff = + cgh.set_specialization_constant(get_coefficients()); cgh.parallel_for( in.get_range(), [=](cl::sycl::item<2> item_id) { From 6510a7617c6d9375c752a6dce1ec120f6fbdd7fe Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Tue, 7 Jan 2020 14:58:56 +0000 Subject: [PATCH 7/8] Use auto in the code example --- proposals/sycl_modules.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/proposals/sycl_modules.md b/proposals/sycl_modules.md index 7c1c8a6..1f6e83f 100644 --- a/proposals/sycl_modules.md +++ b/proposals/sycl_modules.md @@ -960,8 +960,8 @@ void do_conv(buffer in, buffer out) { auto in_acc = in.get_access(cgh); auto out_acc = out.get_access(cgh); - // Create a specialization constant placeholder. - specialization_constant coeff = + // Create a specialization constant placeholder (sycl::specialization_constant object). + auto coeff = cgh.set_specialization_constant(get_coefficients()); cgh.parallel_for( From eec1a5944975f945377426503f3d3dbd82066761 Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Mon, 2 Mar 2020 18:20:58 +0000 Subject: [PATCH 8/8] Rewrite proposal using the suggested kernel_handler. This new proposal remove the need to use `spec_constant` object and uses a kernel handler instead to retrieve the values. Signed-off-by: Victor Lomuller --- proposals/sycl_kernel_handler.md | 94 +++++++++ proposals/sycl_modules.md | 342 +++++++++++++++---------------- 2 files changed, 260 insertions(+), 176 deletions(-) create mode 100644 proposals/sycl_kernel_handler.md diff --git a/proposals/sycl_kernel_handler.md b/proposals/sycl_kernel_handler.md new file mode 100644 index 0000000..7b9ea90 --- /dev/null +++ b/proposals/sycl_kernel_handler.md @@ -0,0 +1,94 @@ +# SYCL kernel_handle + +| | | +| ---------------- | ---------------------------------------| +| Name | kernel_handle | +| Date of creation | 17th Feb 2020 | +| Last updated | 17th Feb 2020 | +| Status | WIP | +| Current revision | 1 | +| Available | N/A | +| Reply-to | Victor Lomuller | +| Original author | Victor Lomuller | +| Contributors | TBD | + +## Overview + +The library implementation of certain device features require the use of a device side handler. +This proposal introduces a `kernel_handler` to provide access to extra device capabilities implementable as a library. + +## Motivation + +The initial proposal of specialization constant forced the user to explicitly get individual `specialization_constant` objects that needed to be propagated in the final program. + +Using a handler, the user has only 1 object to carry to access `specialization_constant` objects thus simplifying the interface. + +For now, the proposal is limited to support specialization constant but can be extended to handle barriers or other functionalities. + +## Revisions + +v1: + + * Initial proposal + +## `sycl::kernel_handler` + +The `sycl::kernel_handler` is a non-user constructible class only passed to user as an argument of the functor passed to `handler::parallel_for` and `handler::parallel_for_work_group`. + +```cpp +namespace sycl { +class kernel_handler { +private: + kernel_handler(); + +public: + + // Return the value associate with the specialization constant id `s`. + // The value returned is either the + template& s> + T get_specialization_constant(); + + template + typename std::remove_reference_t::type get_specialization_constant(); + +}; +} +``` + +## Update `sycl::handler` class definition + +Functor passed to `sycl::handler::single_task`, `sycl::handler::parallel_for` and `sycl::handler::parallel_for_work_group` can take an extra `sycl::kernel_handler` as extra by-value argument. + +Below is an example of invoking a SYCL kernel function with `single_task`: + +```cpp +myQueue.submit([&](handler & cgh) { + cgh.single_task([=] () {}); +}); +``` + +or + +```cpp +myQueue.submit([&](handler & cgh) { + cgh.single_task([=] (sycl::kernel_handler h) {}); +}); +``` + +Below is an example of invoking a SYCL kernel function with `parallel_for`: + +```cpp +myQueue.submit([&](handler & cgh) { + cgh.parallel_for(range<1>(numWorkItems), + [=] (id<1> index) {}); +}); +``` + +or + +```cpp +myQueue.submit([&](handler & cgh) { + cgh.parallel_for(range<1>(numWorkItems), + [=] (id<1> index, sycl::kernel_handler h) {}); +}); +``` diff --git a/proposals/sycl_modules.md b/proposals/sycl_modules.md index 1f6e83f..d87b7c1 100644 --- a/proposals/sycl_modules.md +++ b/proposals/sycl_modules.md @@ -106,10 +106,15 @@ in the scope of the current wording for this proposal to address the ## Revisions +v 0.11: + +* Add a `kernel_hanlder` object to handle access to specialization constant values +* Remove `specialization_constant` in favor of `kernel_hanlder` member function + v 0.10: * Rework of the specialization constant proposal - * Add `specialization _id` and `specialization _constant` definitions + * Add `specialization _id` and `specialization_constant` definitions * Add API to manipulate a Module specialization constants v 0.9: @@ -311,7 +316,7 @@ A *SYCL module* objects may expose *specialization constants*. Specialization is intended for constant objects that will not have known constant values until after initial generation of a module in an intermediate representation format (e.g. SPIR-V). When a module contains *specialization constants*, the method -*`sycl::module::has_specialization_constant()`* returns `true`. +*`sycl::module::use_specialization_constant()`* returns `true`. Note that, although a module may expose *specialization constants*, not all *device images* inside the module may natively support them, the method *`sycl::module::native_specialization_constant()`* returns `true` if the module is capable of using *specialization constants* as immediate value when compiling the module image. When *specialization constants* are available but the user doesn't set values on them, the default values are used. @@ -381,7 +386,7 @@ class module { * Returns true if the current module uses specialization constants. * */ - bool has_specialization_constant() const noexcept; + bool use_specialization_constant() const noexcept; /** * Returns true if the current module can support specialization constants natively. @@ -389,19 +394,15 @@ class module { */ bool native_specialization_constant() const noexcept; - // Available if module_status::executable // Retrieve the value of the specialization constant associated with this module. - // Call only valid from the host. template& s> - specialization_constant get_specialization_constant(handler&); + T get_specialization_constant() const; - // Available if module_status::executable - // Only from C++17 + // Only from C++17 // Set the value of the specialization constant. - // Call only valid from the host. template - specialization_constant::type, s> - get_specialization_constant(handler&); + typename std::remove_reference_t::type + get_specialization_constant() const; // Available if module_status::input // Set the value of the specialization constant. @@ -415,6 +416,15 @@ class module { // Call only valid from the host. template void set_specialization_constant(typename std::remove_reference_t::type); + + // Return true if the specialization constant s is known to the module. + template& s> + bool has_specialization_constant() const noexcept; + + // Only from C++17 + // Return true if the specialization constant s is known to the module. + template + bool has_specialization_constant() const noexcept; }; } // namespace sycl @@ -734,12 +744,66 @@ Specialization constants are associated with the SYCL 1.2.1 program class, and named using a C++ type like kernels. The program class gains a `set_specialization_constant` method, that sets a runtime value to the specialization constant. -The program can then be build with `build_with_kernel_type` to create a +The program can then be built with `build_with_kernel_type` to create a specialized kernel. +### Example + +The following code is showcasing how specialization constant can be used with modules. +The code performs a simple convolution using a scaled convolution kernel. +The specialization constants are here used to inject, at runtime, +the coefficients of the kernel. + +```cpp +#include + +using namespace sycl; + +using coeff_t = std::array, 3>; + +// Read coefficients from somewhere. +coeff_t get_coefficients(); + +// Identify the specialization constant. +specialization_id coeff_id; + +void do_conv(buffer in, buffer out) { + queue myQueue; + + myQueue.submit([&](handler &cgh) { + auto in_acc = in.get_access(cgh); + auto out_acc = out.get_access(cgh); + + // Set the coefficient of the convolution as constant. + // This will build a specific kernel the coefficient available as literals. + cgh.set_specialization_constant(get_coefficients()); + + cgh.parallel_for( + in.get_range(), [=](item<2> item_id, kernel_handler h) { + float acc = 0; + coeff_t coeff = h.get_specialization_constant(); + for (int i = -1; i <= 1; i++) { + if (item_id[0] + i < 0 || item_id[0] + i >= in_acc.get_range()[0]) + continue; + for (int j = -1; j <= 1; j++) { + if (item_id[1] + j < 0 || item_id[1] + j >= in_acc.get_range()[1]) + continue; + // the underlying JIT can see all the values of the array returned by coeff.get(). + acc += coeff[i + 1][j + 1] * + in_acc[item_id[0] + i][item_id[1] + j]; + } + } + out_acc[item_id] = acc; + }); + }); + + myQueue.wait(); +} +``` + ### Specialization constant types -On the new, type-safe, module approach, specialization constants are +In the new, type-safe, module approach, specialization constants are typically associated with SPIR-V module images but are not limited to such images. The specialization constant type is defined as a general SYCL type as below: @@ -765,111 +829,104 @@ public: explicit constexpr specialization_id(Args&&...); }; -template & s> -class specialization_constant { -public: - using type = T; - - // Create an empty (invalid) specialization constant. - // Remains invalid until initialized by a call - // to handler::set_specialization_constant or - // to handler::get_specialization_constant - specialization_constant(); - - specialization_constant(const specialization_constant&) = default; - specialization_constant& operator=(const specialization_constant&) = default; - - T get() const; // explicit access. - operator T() const; // implicit conversion. -}; - } // namespace sycl ``` -`specialization_id` is used by the *SYCL implementation* to *reference* constant whose -concrete value will only be known during the application execution. +`specialization_id` objects are used by the *SYCL implementation* to *reference* constant whose concrete value will only be known during the application execution. To create a new `specialization_id` in a module, the user creates a `specialization_id` object with an *automatic* or *static* storage duration in the namespace or class scope. -The user then uses references to this object to bind the `specialization_constant` value to the *identifier*. `specialization_id` objects must be forward declarable and cannot be moved or copied. -`specialization_constant` is a *placeholder type* used by the *SYCL implementation* to -inject the real value of the specialization constant at runtime. -For the module associated with the host device, and possibly any device with non native specialization constant support, the value is carried by the `specialization_constant` object itself. -For modules natively supporting specialization constant, implementations must guaranty the values returned by the usage of the `specialization_constant` materialized into a constant. - -### Setting the value of specialization constant - -To modify the value of a specialization constant, the following API is available. - -Given a SYCL module, a specialization constant can be set by the `sycl::set_specialization_constant` function. +### Querying if a specialization constant is known to a module ```cpp -namespace sycl { - -template& s> -void set_specialization_constant(module &syclModule, T); +template +class module { +public: +[...] + // Return true if the specialization constant s is known to the module. + template& s> + bool has_specialization_constant() const noexcept; -} // namespace sycl + // Only from C++17 + // Return true if the specialization constant s is known to the module. + template + bool has_specialization_constant() const noexcept; +[...] +}; ``` -If C++17 or later is enabled, then the following overload is available +The function `has_specialization_constant` returns true if the `specialization_id` id is known to the current module. +If no values were ever set for the module, but the id is known to the module, the function still returns true. +This function only checks if it is meaningful to get or set the value of a `specialization_id` with a given module. -```cpp -namespace sycl { +### Setting the value of specialization constant of a module -template -void set_specialization_constant(module &syclModule, - typename std::remove_reference_t::type); +A specialization constant for a module can be set while in the `input` state. -} // namespace sycl +```cpp +template +class module { +public: +[...] + template& s> + void set_specialization_constant(T); + + // Available if module_status::input + // Only from C++17 + // Set the value of the specialization constant. + // Call only valid from the host. + template + void set_specialization_constant(typename std::remove_reference_t::type); +[...] +}; ``` -It is valid for a user to set several times a specialization constant with the same `specialization_id`, -only the last version of the value will be consider when building the module. -Calling this function is not allowed during the submission of a task. +The function `set_specialization_constant` binds the runtime value of type `T` to the module for the `specialization_id` identifier. +If a value was already set for the given id then the value is overwritten by the new value. -### Getting the value of specialization constant in kernel +For all versions of `set_specialization_constant`, if the `specialization_id` is unknown to the module (`has_specialization_constant` returns false for this id) then a runtime exception is thrown. -To retrieve the value of a specialization constant in a SYCL kernel context, the following API is available. +### Getting the value of a specialization constant of a module ```cpp -namespace sycl { - -template& s> -T get_specialization_constant(const specialization_constant&); +template +class module { +public: +[...] + template& s> + T get_specialization_constant() const; -} // namespace sycl + // Only from C++17 + template + typename std::remove_reference_t::type + get_specialization_constant() const; +[...] +}; ``` -If C++17 or later is enabled, then the following overload is available -```cpp -namespace sycl { -template -typename std::remove_reference_t::type -get_specialization_constant(const specialization_constant::type, s>&); +The `get_specialization_constant` returns the value identified by `specialization_id` that is bound to this module. +If no value was set for the id, then the function returns the default value associated with the id. -} // namespace sycl -``` +For all versions of `get_specialization_constant`, if the `specialization_id` is unknown to the module (`has_specialization_constant` returns false for this id) then a runtime exception is thrown. + +### Getting the value of specialization constant in kernel -This call is only valid whitin a SYCL kernel context and the `specialization_constant` value must have been -created by the handler that invoked the kernel from which the `sycl::get_specialization_constant` is called. +Values of specialization constants can be retrieved inside kernels using the `kernel_handler` function `get_specialization_constant`. -If the module has native support for specialization constant, -the returned value becomes a constant value available for the runtime compiler to use for optimisations. +A call to `kernel_handler::get_specialization_constant` will return the value bound to the id and the module which has been used to run the kernel/ +If the module has native support for specialization constant, the returned value becomes a literal value available for the runtime compiler to use for optimizations. +Calls to `kernel_handler::get_specialization_constant` are constant and never throw. **Note:** the value may become a constant (known numerical value), although this may only happen at runtime. From the SYCL compiler's perspective (C++ compiler) the value is understood to be a runtime value and -for this reason it is not possible to use such value in a `constexpr`. +for this reason, it is not possible to use such value as a `constexpr` value. -### Retrieving a specialization constant placeholder +Usage of a `specialization_id` via the `kernel_handler` makes the device compiler bind the id to the module in which the will kernel lives, making it "known" to the module. -A `specialization_constant` handler can be retrieved from the command group handler by either: +### Setting and getting a specialization constant during a submit -- Getting the placeholder for a previously set specialization constant value; - - If no specialization constant where ever set, then the *placeholder* object holds the default value of its `specialization_id`. -- Setting a specialization constant value and retrieving the placeholder - - Performance notice: this function may cause the SYCL runtime to implicitly recompile the relevant SYCL module with the given value. This may have an important effect on performances and users should check their vendor recommendations on how to efficiently uses this functionality. +Specialization constant values can be set/retrieved directly via the handler object. The command group `handler` class is extended with the following API: @@ -879,109 +936,42 @@ namespace sycl { // ... public: - template& s> - specialization_constant set_specialization_constant(T); - - template& s> - specialization_constant get_specialization_constant(); -// ... -}; - -} // namespace sycl -``` - -If C++17 or later is enabled, then the following overloads are available - -```cpp -namespace sycl { - class handler { -// ... -public: - - // Same as template& s> specialization_constant set_specialization_constant(T) + // Set the value of a specialization constant prior to invocation. + template& s> + void set_specialization_constant(T); + // From C++17 + // Set the value of a specialization constant prior to invocation. template - specialization_constant::type, s> - set_specialization_constant(typename std::remove_reference_t::type); + void set_specialization_constant(typename std::remove_reference_t::type); - // Same as template& s> specialization_constant get_specialization_constant() + // Get the value of a specialization constant prior to invocation. + template& s> + T get_specialization_constant() const; + // From C++17 + // Get the value of a specialization constant prior to invocation. template - specialization_constant::type, s> - get_specialization_constant(); + typename std::remove_reference_t::type get_specialization_constant() const; // ... }; } // namespace sycl ``` -| Member function | Description | -|-----------------------------------------------------------------------------|-------------------------------------------------------------------------------------------------------------------------------------------| -| `template& s> specialization_constant set_specialization_constant(T)` | Set the given value as the specialization constant value in the underlying module and returns a `specialization_constant` handler usable in kernel. If the user specify to the handle a specific module to use, then the value given to `set_specialization_constant` must match the the one used to build the module. In case of mismatch, the SYCL runtime raises an error. | -| `template& s> specialization_constant get_specialization_constant()` | Returns a `specialization_constant` handler usable in kernel. | - -Only one call to `set_specialization_constant` or `get_specialization_constant` is allowed per `specialization_id` and per command group scope. -If more than one call per `specialization_id` and per command group scope is made then an runtime exception is thrown. +The function `set_specialization_constant` binds the runtime value of type `T` to underlying module for the `specialization_id` identifier. +If the `specialization_id` is unknown to the underlying module a runtime exception is thrown. +If a module were specified (using `handler::use_module`), the value given to `set_specialization_constant` must be consistent with the value used to build the specified module. +If the values are not consistent then a runtime exception is thrown. +*Implementation note*: the use of `set_specialization_constant` can force the runtime to rebuild a module if one of the specialization constant value changes. +This may impact performance. Users are invited to check the documentation of their implementation on how to optimally use this function. -Note the value of the specialization constant depends on the module that is used, -not on the placeholder object. +The `get_specialization_constant` returns the value identified by `specialization_id` that is bound to the underlying module. +If no value was set for the id, then the function returns the default value associated with the id. -*Implementation note*: Although specialization constants are meant for situations -where there is a final online compilation stage, if the implementation supports -compiling directly on the backend for an ISA, specialization constants can still -be supported via two mechanism: (1) Having different specialized images for known +*Implementation note*: specialization constants are meant for situations +where there is a final online compilation stage. Although if the implementation statically +compiles directly to an ISA, specialization constants can still +be supported via two mechanisms: (1) Having different specialized images for known values on the specialization constant or (2) having the value of the specialization constant as a normal kernel parameter. If choosing (1), the original SPIR-V image should be part of the SYCL module so that the module can be specialized for non-known values. - -### Example - -The following code is showcasing how specialization constant can be used. -The code perform a simple convolution using a scaled convolution kernel. -The specialization constants are here used to inject, at runtime, -the coefficients of the kernel. - -```cpp -#include - -using namespace cl::sycl; - -using coeff_t = std::array, 3>; - -// Read coefficients from somewhere. -coeff_t get_coefficients(); - -// Identify the specialization constant. -specialization_id coeff_id; - -void do_conv(buffer in, buffer out) { - queue myQueue; - - myQueue.submit([&](handler &cgh) { - auto in_acc = in.get_access(cgh); - auto out_acc = out.get_access(cgh); - - // Create a specialization constant placeholder (sycl::specialization_constant object). - auto coeff = - cgh.set_specialization_constant(get_coefficients()); - - cgh.parallel_for( - in.get_range(), [=](cl::sycl::item<2> item_id) { - float acc = 0; - for (int i = -1; i <= 1; i++) { - if (item_id[0] + i < 0 || item_id[0] + i >= in_acc.get_range()[0]) - continue; - for (int j = -1; j <= 1; j++) { - if (item_id[1] + j < 0 || item_id[1] + j >= in_acc.get_range()[1]) - continue; - // the underlying JIT can see all the values of the array returned by coeff.get(). - acc += coeff.get()[i + 1][j + 1] * - in_acc[item_id[0] + i][item_id[1] + j]; - } - } - out_acc[item_id] = acc; - }); - }); - - myQueue.wait(); -} -```