From d23be44e9f4bacff0636664d81a898ed63b3d220 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Fri, 1 Mar 2024 13:28:56 -0500 Subject: [PATCH 01/10] Minor formatting fixes --- ...ext_oneapi_kernel_compiler_opencl.asciidoc | 32 ++++++++------- ..._ext_oneapi_kernel_compiler_spirv.asciidoc | 39 +++++++++++-------- 2 files changed, 40 insertions(+), 31 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc index 32ebd6b2a2ba6..1f6aa202b5771 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc @@ -21,7 +21,7 @@ == Notice [%hardbreaks] -Copyright (C) 2023-2023 Intel Corporation. All rights reserved. +Copyright (C) 2023-2024 Intel Corporation. All rights reserved. Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks of The Khronos Group Inc. @@ -49,11 +49,11 @@ This extension also depends on the following other SYCL extensions: == 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 is an experimental extension specification, intended to provide early +access to features and gather community feedback. Interfaces defined in +this specification are implemented in {dpcpp}, but they are not finalized +and may change incompatibly in future versions of {dpcpp} without prior notice. +*Shipping software products should not rely on APIs defined in this specification.* @@ -96,7 +96,8 @@ This extension adds the `opencl` enumerator to the `source_language` enumeration, which indicates that a kernel bundle defines kernels in the OpenCL C language. -``` +[source,c++] +---- namespace sycl::ext::oneapi::experimental { enum class source_language : /*unspecified*/ { @@ -105,7 +106,7 @@ enum class source_language : /*unspecified*/ { }; } // namespace sycl::ext::oneapi::experimental -``` +---- === Source code is text format @@ -132,17 +133,19 @@ functions identify a kernel using the function name, exactly as it appears in the OpenCL C source code. For example, if the kernel is defined this way in OpenCL C: -``` +[source,c++] +---- __kernel void foo(__global int *in, __global int *out) {/*...*/} -``` +---- Then the application's host code can query for the kernel like so: -``` +[source,c++] +---- sycl::kernel_bundle kb = /*...*/; sycl::kernel k = kb.ext_oneapi_get_kernel("foo"); -``` +---- === Kernel argument restrictions @@ -192,7 +195,8 @@ _{endnote}_] The following example shows a simple SYCL program that defines an OpenCL C kernel as a string and then compiles and launches it. -``` +[source,c++] +---- #include #include namespace syclex = sycl::ext::oneapi::experimental; @@ -241,7 +245,7 @@ int main() { cgh.parallel_for(sycl::range{N}, k); }); } -``` +---- == Issues diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc index c39f1fd6ae37a..f5156968e9f13 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc @@ -10,6 +10,7 @@ :encoding: utf-8 :lang: en :dpcpp: pass:[DPC++] +:cpp: pass:[C++] :endnote: —{nbsp}end{nbsp}note // Set the default source code type in this document to C++, @@ -21,7 +22,7 @@ == Notice [%hardbreaks] -Copyright (C) 2023-2023 Intel Corporation. All rights reserved. +Copyright (C) 2023-2024 Intel Corporation. All rights reserved. Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks of The Khronos Group Inc. @@ -52,11 +53,11 @@ This extension also depends on the following other SYCL extensions: == 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 is an experimental extension specification, intended to provide early +access to features and gather community feedback. Interfaces defined in +this specification are implemented in {dpcpp}, but they are not finalized +and may change incompatibly in future versions of {dpcpp} without prior notice. +*Shipping software products should not rely on APIs defined in this specification.* @@ -103,7 +104,8 @@ This extension adds the `spirv` enumerator to the `source_language` enumeration, which indicates that a kernel bundle defines kernels as a SPIR-V binary module. -``` +[source,c++] +---- namespace sycl::ext::oneapi::experimental { enum class source_language : /*unspecified*/ { @@ -112,7 +114,7 @@ enum class source_language : /*unspecified*/ { }; } // namespace sycl::ext::oneapi::experimental -``` +---- === Source code is binary format @@ -155,16 +157,18 @@ identify a kernel using the name, exactly as it appears in the *OpEntryPoint* instruction. For example, if the kernel is defined this way in SPIR-V: -``` +[source] +---- OpEntryPoint Kernel %20 "foo" -``` +---- Then the application's host code can query for the kernel like so: -``` +[source,c++] +---- sycl::kernel_bundle kb = /*...*/; sycl::kernel k = kb.ext_oneapi_get_kernel("foo"); -``` +---- === Passing kernel arguments @@ -182,7 +186,7 @@ backend's SPIR-V client API specification. |Corresponding SYCL type |*OpTypeInt* -|A C++ type that is device copyable, which has the same width and data +|A {cpp} type that is device copyable, which has the same width and data representation. [_Note:_ Applications typically use a fixed-width integer type where the width @@ -190,7 +194,7 @@ matches the width of the *OpTypeInt*. _{endnote}_] |*OpTypeFloat* -|A C++ type that is device copyable, which has the same width and data +|A {cpp} type that is device copyable, which has the same width and data representation. [_Note:_ Applications typically use `float` when the *OpTypeFloat* has a width @@ -207,7 +211,7 @@ _{endnote}_] |*OpTypePointer* with _Storage Class_ *Function* and _Type_ *OpTypeStruct* (i.e. the pointed-at type is *OpTypeStruct*). -|A C++ struct or class that is device copyable, which has the same size and +|A {cpp} struct or class that is device copyable, which has the same size and data representation as the *OpTypeStruct*. [_Note:_ The SYCL argument is a structure even though the SPIR-V argument type @@ -236,7 +240,8 @@ _{endnote}_] The following example shows a simple SYCL program that loads a SPIR-V module from a file and then launches a kernel from that module. -``` +[source,c++] +---- #include #include #include @@ -288,4 +293,4 @@ int main() { cgh.parallel_for(sycl::range{N}, k); }); } -``` +---- From 320a65ecdc289e5ebc4cc5453afc19af926b30db Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Fri, 1 Mar 2024 17:56:43 -0500 Subject: [PATCH 02/10] Fix name of OpenCL header in example --- .../sycl_ext_oneapi_kernel_compiler_opencl.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc index 1f6aa202b5771..1468be127ac27 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc @@ -198,7 +198,7 @@ kernel as a string and then compiles and launches it. [source,c++] ---- #include -#include +#include namespace syclex = sycl::ext::oneapi::experimental; int main() { From d9afed6a64e60dc862eb78ee59bf3d9b6304b506 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Fri, 1 Mar 2024 13:26:53 -0500 Subject: [PATCH 03/10] Clarify local accessor arguments to SPIR-V --- .../sycl_ext_oneapi_kernel_compiler_spirv.asciidoc | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc index f5156968e9f13..f7859b47d2292 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc @@ -209,6 +209,14 @@ _{endnote}_] |*OpTypePointer* with _Storage Class_ *Workgroup* |A `local_accessor`. +[_Note:_ The `local_accessor` merely conveys the size of the local memory, such +that the *OpTypePointer* kernel argument points to a local memory buffer of _N_ +bytes, where _N_ is the value returned by `local_accessor::byte_size`. +If the application wants to pass other information from the `local_accessor` to +the kernel (such as the value _N_), it must pass this as separate kernel +arguments. +_{endnote}_] + |*OpTypePointer* with _Storage Class_ *Function* and _Type_ *OpTypeStruct* (i.e. the pointed-at type is *OpTypeStruct*). |A {cpp} struct or class that is device copyable, which has the same size and From 99a06fd837e4f6e2957c5674a3aaadcd7aa2954f Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Fri, 1 Mar 2024 15:40:33 -0500 Subject: [PATCH 04/10] Clarify iteration space and builtin vars in SPIR-V Also disallows launching a SPIR-V kernel as a simple "range" kernel. This seems consistent with our view that a "range" kernel is not just a degenerate form of an nd-range kernel. Since SPIR-V kernels always have access to nd-range features, it does not make sense to launch them as range kernels. This is also consistent with our decision to limit SYCL free function kernels to "nd-range" and "single-task" forms. --- ..._ext_oneapi_kernel_compiler_spirv.asciidoc | 34 +++++++++++++++++-- 1 file changed, 32 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc index f7859b47d2292..858031c72c1a8 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc @@ -242,6 +242,34 @@ because this helps ensure that the size of the integers on the host matches the size in the kernel. _{endnote}_] +=== Iteration space and built-in variables + +A `kernel` object created from a SPIR-V module must be launched either as a +single-task kernel or as an nd-range kernel. +Attempting to launch such a kernel with a simple range iteration space results +in undefined behavior. + +If the kernel is launched as a single-task kernel, it is executed with one +work-group dimension, with one work-group of one work-item. + +If the kernel is launched as an nd-range kernel, the number of work-group +dimensions is the same as the number of dimensions in the `nd_range`. +The global size, local size, and the number of work-groups is determined in the +usual way from the `nd_range`. +If the SPIR-V kernel defines a fixed local size (e.g. via the *LocalSize* +execution mode), the local size in the `nd_range` must match this value. + +The kernel may use SPIR-V built-in variables as they are defined in the SPIR-V +client API specification, with the following clarification. +Some of the built-in variables are defined as a 3-component vector (e.g. +*GlobalSize*). +The order of these components is inverted from their corresponding SYCL +objects. +To illustrate, consider a 3-dimensional kernel invocation. +Vector component 0 of the *GlobalSize* variable corresponds to dimension 2 of +the `nd_range`, and vector component 2 of the *GlobalSize* variable corresponds +to dimension 0 of the `nd_range`. + == Example @@ -284,6 +312,7 @@ int main() { sycl::kernel k = kb_exe.ext_oneapi_get_kernel("my_kernel"); constexpr int N = 4; + constexpr int WGSIZE = 1; int32_t input[N] = {0, 1, 2, 3}; int32_t output[N] = {}; @@ -297,8 +326,9 @@ int main() { // Set the values for the kernel arguments. cgh.set_args(in, out); - // Invoke the kernel over a range. - cgh.parallel_for(sycl::range{N}, k); + // Invoke the kernel over an nd-range. + sycl::nd_range ndr{{N}, {WGSIZE}}; + cgh.parallel_for(ndr, k); }); } ---- From 8433e394652f96e8187344c21906ef20c524a0a9 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Fri, 1 Mar 2024 16:49:01 -0500 Subject: [PATCH 05/10] Align OpenCL kernel arg table with SPIR-V The OpenCL table describing legal kernel arguments had the opposite order from the corresponding SPIR-V table. The OpenCL table listed the SYCL argument type in the left column while the SPIR-V table listed the SYCL argument tupe in the right column. I like the SPIR-V format better, so change the OpenCL table to match. This also adopts some of the same wording from the SPIR-V table into the OpenCL one. --- ...ext_oneapi_kernel_compiler_opencl.asciidoc | 64 +++++++++++-------- 1 file changed, 36 insertions(+), 28 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc index 1468be127ac27..c4a1277f1f473 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc @@ -149,46 +149,54 @@ sycl::kernel k = kb.ext_oneapi_get_kernel("foo"); === Kernel argument restrictions -When a kernel is defined in OpenCL C and invoked from SYCL via a `kernel` -object, the arguments to the kernel are restricted to certain types. -In general, the host application passes an argument value via -`handler::set_arg` using one type and the kernel receives the argument value -as a corresponding OpenCL C type. -The following table lists the set of valid types for these kernel arguments: - +The following table defines the set of OpenCL C kernel argument types that are +supported by this extension and explains how to pass each type of argument from +SYCL. [%header,cols="1,1"] |=== -|Type in SYCL host code -|Type in OpenCL C kernel +|OpenCL C type +|Corresponding SYCL type -|One of the OpenCL scalar types (e.g. `cl_int`, `cl_float`, etc.) -|The corresponding OpenCL C type (e.g. `int`, `float`, etc.) +|One of the OpenCL C scalar types (e.g. `int`, `float`, etc.) +|A {cpp} type that is device copyable, which has the same width and data + representation. -|A USM pointer. -|A `+__global+` pointer of the corresponding type. +[_Note:_ Applications typically use the corresponding OpenCL type (e.g. +`cl_int`, `cl_float`, etc.) +_{endnote}_] -|A class (or struct) that is device copyable in SYCL whose elements are - composed of OpenCL scalar types or USM pointers. -|A class (or struct) passed by value whose elements have the corresponding - OpenCL C types. +|A `+__global+` pointer. +|Either a {cpp} pointer (typically a pointer to USM memory) or an `accessor` + whose target is `target::device`. -|An `accessor` with `target::device` whose `DataT` is an OpenCL scalar type, - a USM pointer, or a device copyable class (or struct) whose elements are - composed of these types. -|A `+__global+` pointer to the first element of the accessor's buffer. - The pointer has the corresponding OpenCL C type. +|A `+__local+` pointer. +|A `local_accessor`. -[_Note:_ The accessor's size is not passed as a kernel argument, so the host -code must pass a separate argument with the size if this is desired. +[_Note:_ The `local_accessor` merely conveys the size of the local memory, such +that the kernel argument points to a local memory buffer of _N_ bytes, where +_N_ is the value returned by `local_accessor::byte_size`. +If the application wants to pass other information from the `local_accessor` to +the kernel (such as the value _N_), it must pass this as separate kernel +arguments. _{endnote}_] -|A `local_accessor` whose `DataT` is an OpenCL scalar type, a USM pointer, or a - device copyable class (or struct) whose elements are composed of these types. -|A `+__local+` pointer to the first element of the accessor's local memory. - The pointer has the corresponding OpenCL C type. +|A class (or struct) passed by value. +|A {cpp} struct or class that is device copyable, which has the same size and + data representation as the OpenCL C struct. + +[_Note:_ The SYCL argument must not contain any `accessor` or `local_accessor` +members because these types are not device copyable. +If the OpenCL C structure contains a pointer member, the corresponding SYCL +structure member is typically a USM pointer. +_{endnote}_] |=== +When data allocated on the host is accessed by the kernel via a pointer, the +application must ensure that the data has the same size and representation on +the host and inside the OpenCL C kernel. +Applications can use the OpenCL types (e.g. `cl_int`) for this purpose. + == Example From ae01c9e33e6aa335e96bf5949304601cb3f068f9 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Fri, 1 Mar 2024 17:16:18 -0500 Subject: [PATCH 06/10] Clarify iteration space and work-item functions Clarify the iteration space and the user of work-item functions in the OpenCL spec. These changes closely follow the corresponding changes in the SPIR-V spec. --- ...ext_oneapi_kernel_compiler_opencl.asciidoc | 51 ++++++++++++------- 1 file changed, 32 insertions(+), 19 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc index c4a1277f1f473..9241aea64cf34 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc @@ -197,6 +197,34 @@ application must ensure that the data has the same size and representation on the host and inside the OpenCL C kernel. Applications can use the OpenCL types (e.g. `cl_int`) for this purpose. +=== Iteration space and work-item functions + +A `kernel` object created from OpenCL C source code must be launched either as +a single-task kernel or as an nd-range kernel. +Attempting to launch such a kernel with a simple range iteration space results +in undefined behavior. + +If the kernel is launched as a single-task kernel, it is executed with one +work-group dimension, with one work-group of one work-item. + +If the kernel is launched as an nd-range kernel, the number of work-group +dimensions is the same as the number of dimensions in the `nd_range`. +The global size, local size, and the number of work-groups is determined in the +usual way from the `nd_range`. +If the OpenCL C kernel is decorated with the `reqd_work_group_size` attribute, +the local size in the `nd_range` must match this value. + +The kernel may call the functions defined in section 6.15.1 "Work-Item +Functions" of the OpenCL C specification, with the following clarification. +Some of these functions take a `dimindx` parameter that selects a dimension +index. +This index has the opposite sense from SYCL. +To illustrate, consider a call to `get_global_size` from a kernel that is +invoked with a 3-dimensional `nd_range`. +Calling `get_global_size(0)` retrieves the global size from dimension 2 of the +`nd_range`, and calling `get_global_size(2)` retrieves the global size from +dimension 0 of the `nd_range`. + == Example @@ -236,6 +264,7 @@ int main() { sycl::kernel k = kb_exe.ext_oneapi_get_kernel("my_kernel"); constexpr int N = 4; + constexpr int WGSIZE = 1; cl_int input[N] = {0, 1, 2, 3}; cl_int output[N] = {}; @@ -249,8 +278,9 @@ int main() { // Each argument to the kernel is a SYCL accessor. cgh.set_args(in, out); - // Invoke the kernel over a range. - cgh.parallel_for(sycl::range{N}, k); + // Invoke the kernel over an nd-range. + sycl::nd_range ndr{{N}, {WGSIZE}}; + cgh.parallel_for(ndr, k); }); } ---- @@ -301,20 +331,3 @@ their source string. backend? Currently, the online_compiler does support this case (but it provides no way to query about optional features or extensions). - -* Do we need to document some restrictions on the OpenCL C - https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_C.html#work-item-functions[ - work-item functions] that the kernel can call, which depends on how the - kernel was launched? - For example, can a kernel launched with the simple `range` form of - `parallel_for` call `get_local_size`? - In OpenCL, there is only one way to launch kernels - (`clEnqueueNDRangeKernel`), so it is always legal to call any of the - work-item functions. - If an OpenCL kernel is launched with a NULL `local_work_size` (which is - roughly equivalent to SYCL's `range` form of `parallel_for`), the - `get_local_size` function returns the local work-group size that is chosen by - the implementation. - Level Zero, similarly, has only one way to launch kernels. - Therefore, maybe it is OK to let kernels in this extension call any of the - work-item functions, regardless of how they are launched? From dd9fb988a78bab9bdacc05d37842c48aadea5992 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Tue, 5 Mar 2024 15:01:17 -0500 Subject: [PATCH 07/10] Improve wording of single-task kernels --- .../sycl_ext_oneapi_kernel_compiler_opencl.asciidoc | 4 ++-- .../sycl_ext_oneapi_kernel_compiler_spirv.asciidoc | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc index 9241aea64cf34..0e8aadf1f9290 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc @@ -204,8 +204,8 @@ a single-task kernel or as an nd-range kernel. Attempting to launch such a kernel with a simple range iteration space results in undefined behavior. -If the kernel is launched as a single-task kernel, it is executed with one -work-group dimension, with one work-group of one work-item. +If the kernel is launched as a single-task kernel, it is executed with a +1-dimensional nd-range, with one work-group of one work-item. If the kernel is launched as an nd-range kernel, the number of work-group dimensions is the same as the number of dimensions in the `nd_range`. diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc index 858031c72c1a8..6c0b14a97b35a 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc @@ -249,8 +249,8 @@ single-task kernel or as an nd-range kernel. Attempting to launch such a kernel with a simple range iteration space results in undefined behavior. -If the kernel is launched as a single-task kernel, it is executed with one -work-group dimension, with one work-group of one work-item. +If the kernel is launched as a single-task kernel, it is executed with a +1-dimensional nd-range, with one work-group of one work-item. If the kernel is launched as an nd-range kernel, the number of work-group dimensions is the same as the number of dimensions in the `nd_range`. From 3960d76dee0f74054301b0c0c09375390183e01c Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Wed, 6 Mar 2024 14:54:54 -0500 Subject: [PATCH 08/10] Reference C.7.7 when talking about index inversion Add a note referencing section C.7.7 of the core SYCL spec, which has more details about the inverted sense of SYCL dimensions compared to OpenCL. --- .../sycl_ext_oneapi_kernel_compiler_opencl.asciidoc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc index 0e8aadf1f9290..70149da66a5cb 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc @@ -218,7 +218,8 @@ The kernel may call the functions defined in section 6.15.1 "Work-Item Functions" of the OpenCL C specification, with the following clarification. Some of these functions take a `dimindx` parameter that selects a dimension index. -This index has the opposite sense from SYCL. +This index has the opposite sense from SYCL, as described in section C.7.7 +"OpenCL kernel conventions and SYCL" of the core SYCL specification. To illustrate, consider a call to `get_global_size` from a kernel that is invoked with a 3-dimensional `nd_range`. Calling `get_global_size(0)` retrieves the global size from dimension 2 of the From 12baf159b510243a9861b1a750079a8692b9ec7f Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Wed, 6 Mar 2024 15:12:11 -0500 Subject: [PATCH 09/10] Clarify single-task behavior Clarify that a single-task OpenCL and SPIR-V kernel can still use features like `local_accessor` and group functions. These are well-defined because we already state that these single-task kernels are launched as nd-range kernels with 1 work-item. --- .../sycl_ext_oneapi_kernel_compiler_opencl.asciidoc | 7 +++++++ .../sycl_ext_oneapi_kernel_compiler_spirv.asciidoc | 6 ++++++ 2 files changed, 13 insertions(+) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc index 70149da66a5cb..7decd56796722 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc @@ -206,6 +206,13 @@ in undefined behavior. If the kernel is launched as a single-task kernel, it is executed with a 1-dimensional nd-range, with one work-group of one work-item. +Because it is launched as an nd-range kernel, the kernel can use features that +are normally prohibited in single-task kernels. +For example, the `local_accessor` type is allowed as a kernel argument, and the +kernel can use OpenCL C work-group collective functions and sub-group +functions. +Of course, these features have limited use because the kernel is launched with +just a single work-item. If the kernel is launched as an nd-range kernel, the number of work-group dimensions is the same as the number of dimensions in the `nd_range`. diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc index 6c0b14a97b35a..f539ad7dfea24 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc @@ -251,6 +251,12 @@ in undefined behavior. If the kernel is launched as a single-task kernel, it is executed with a 1-dimensional nd-range, with one work-group of one work-item. +Because it is launched as an nd-range kernel, the kernel can use features that +are normally prohibited in single-task kernels. +For example, the `local_accessor` type is allowed as a kernel argument, and the +kernel can use SPIR-V group and subgroup instructions. +Of course, these features have limited use because the kernel is launched with +just a single work-item. If the kernel is launched as an nd-range kernel, the number of work-group dimensions is the same as the number of dimensions in the `nd_range`. From 51c0f9f069f8cc56d0d60584e8f92578abfa6458 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Thu, 7 Mar 2024 09:36:11 -0500 Subject: [PATCH 10/10] Minor formatting change Use [source] block instead of triple backticks. --- .../sycl_ext_oneapi_kernel_compiler_opencl.asciidoc | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc index 67e83236970ba..9a021ed3a52f3 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc @@ -447,7 +447,8 @@ This example demonstrates how to query the version of OpenCL C that is supported, how to query the supported features, and how to query the supported extensions. -``` +[source,c++] +---- #include #include namespace syclex = sycl::ext::oneapi::experimental; @@ -476,4 +477,4 @@ int main() { std::cout << "Device supports online compilation with the OpenCL full profile\n"; } -``` +----