diff --git a/sycl/doc/extensions/experimental/sycl_ext_intel_maximum_registers.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_intel_maximum_registers.asciidoc new file mode 100644 index 0000000000000..2515c937f314b --- /dev/null +++ b/sycl/doc/extensions/experimental/sycl_ext_intel_maximum_registers.asciidoc @@ -0,0 +1,251 @@ += sycl_ext_intel_maximum_registers + +:source-highlighter: coderay +:coderay-linenums-mode: table + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en +:dpcpp: pass:[DPC++] +:endnote: —{nbsp}end{nbsp}note + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + + +== Notice + +[%hardbreaks] +Copyright (C) 2026 Intel Corporation. All rights reserved. + +Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks +of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by +permission by Khronos. + + +== Contact + +To report problems with this extension, please open a new issue at: + +https://github.com/intel/llvm/issues + + +== Dependencies + +This extension is written against the SYCL 2020 revision 11 specification. All +references below to the "core SYCL specification" or to section numbers in the +SYCL specification refer to that revision. + +This extension also depends on the following other SYCL extensions: + +* link:../experimental/sycl_ext_oneapi_properties.asciidoc[ + sycl_ext_oneapi_properties] +* link:../experimental/sycl_ext_oneapi_kernel_properties.asciidoc[ + sycl_ext_oneapi_kernel_properties] + +== Status + +This is an experimental extension specification, intended to provide early access +to features and gather community feedback. Interfaces defined in this specification +are implemented in DPC++, but they are not finalized and may change incompatibly in +future versions of DPC++ without prior notice. **Shipping software products should not +rely on APIs defined in this specification.** + +== Backend support status + +This extension is not yet implemented. + +== Overview + +This extension adds new kernel properties which allow an application to choose +the size of the general register file (GRF) that the compiler uses when +generating device code for a kernel. + +These properties are intended for advanced performance tuning. +Increasing the GRF size may improve performance for kernels with high register +requirements, but it can also introduce trade-offs that reduce performance. +Applications should generally rely on the compiler's default selection unless +profiling demonstrates a benefit from specifying a different GRF size. + +This extension is similar to link:../experimental/sycl_ext_intel_grf_size.asciidoc[sycl_ext_intel_grf_size] +but the mechanism by which the GRF size specification is communicated from {dpcpp} to the runtime differs. + +== Specification + +=== Feature test macro + +This extension provides a feature-test macro as described in the core SYCL +specification. An implementation supporting this extension must predefine the +macro `SYCL_EXT_INTEL_MAXIMUM_REGISTERS` to one of the values defined in the table +below. Applications can test for the existence of this macro to determine if +the implementation supports this feature, or applications can test the macro's +value to determine which of the extension's features the implementation +supports. + +[%header,cols="1,5"] +|=== +|Value +|Description + +|1 +|The APIs of this experimental extension are not versioned, so the + feature-test macro always has this value. +|=== + +=== Properties + +This extension adds the properties listed below, which applications can use +when defining a kernel to select the kernel's GRF size. +When a kernel is defined without either property, the GRF size is +implementation-defined. +On devices that support two GRF sizes, implementations typically select the +smaller size by default. + +A kernel may be defined with at most one of these GRF-setting properties. + +[_Note_: Using these properties to select a larger GRF size may impose a lower +limit on the kernel's maximum work-group size. +Applications can query the maximum work-group size for the kernel via the +`info::kernel_device_specific::work_group_size` information descriptor, which is +defined in the core SYCL specification. +_{endnote}_] + +''' + +[source,c++] +---- +namespace sycl::ext::intel::experimental { + +struct maximum_registers_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +template +inline constexpr maximum_registers_key::value_t maximum_registers; + +} // namespace sycl::ext::intel::experimental +---- + +Indicates that the kernel should be compiled with a GRF size of `Size`. +The value `Size` must be one of the legal GRF size values for the device to +which the kernel is submitted, as defined by the table below. + +''' + +[source,c++] +---- +namespace sycl::ext::intel::experimental { + +struct maximum_registers_automatic_key { + using value_t = + oneapi::experimental::property_value; +}; + +inline constexpr maximum_registers_automatic_key::value_t maximum_registers_automatic; + +} // namespace sycl::ext::intel::experimental +---- + +Indicates that the compiler should use a heuristic to determine the GRF size for +the kernel. +This property gives the compiler more freedom to chooes a larger GRF size +compared to kernels that are not defined with any GRF property. + +=== Legal GRF sizes for Intel devices + +The following table lists the legal GRF sizes for some Intel devices. + +[%header,cols="1,5"] +|=== +|GPU |Supported Values +| PVC | 128 (small register file), 256 (large register file) +| DG2 | 128 (small register file), 256 (large register file) +| CRI | 128, 256, 512 +|=== + +Providing a GRF size that is not consistent with the supported values results in +undefined behavior. + + +== Examples + +=== Basic usage + +A simple example of using this extension is below. + +[source,c++] +---- +#include +namespace syclex = sycl::ext::oneapi::experimental; +namespace intelex = sycl::ext::intel::experimental; + +struct Kernel1 { + Kernel1() {} + + void operator()(sycl::nd_item<1> it) const { /* ... */ } + + auto get(syclex::properties_tag) const { + return syclex::properties{intelex::maximum_registers<256>}; + } +}; + +struct Kernel2 { + Kernel2() {} + + void operator()(sycl::nd_item<1> it) const { /* ... */ } + + auto get(syclex::properties_tag) const { + return syclex::properties{intelex::maximum_registers_automatic}; + } +}; + +int main() { + sycl::queue q; + + sycl::nd_range ndr{{256},{16}}; + q.parallel_for(ndr, Kernel1{}).wait(); + q.parallel_for(ndr, Kernel2{}).wait(); +} +---- + +=== Usage with a free function kernel + +An example using this extension with free function kernels +(link:../experimental/sycl_ext_oneapi_free_function_kernels.asciidoc[ +sycl_ext_oneapi_free_function_kernels]). + +[source,c++] +---- +#include +namespace syclex = sycl::ext::oneapi::experimental; +namespace intelex = sycl::ext::intel::experimental; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclex::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((intelex::maximum_registers<256>)) +void kernel1() { + /* ... */ +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclex::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((intelex::maximum_registers_automatic)) +void kernel2() { + /* ... */ +} + +int main() { + sycl::queue q; + + sycl::nd_range ndr{{256},{16}}; + syclex::nd_launch(q, ndr, syclex::kernel_function); + syclex::nd_launch(q, ndr, syclex::kernel_function); + q.wait(); +} +----