From 1996e80db3751f98676c17ad36efe95ea3ef3e04 Mon Sep 17 00:00:00 2001 From: Konstantin S Bobrovsky Date: Mon, 15 Aug 2022 13:48:31 -0700 Subject: [PATCH 1/5] [SYCL] Clarify sub-group size calculation in invoke_simd spec. Signed-off-by: Konstantin S Bobrovsky --- .../sycl_ext_oneapi_invoke_simd.asciidoc | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc index fd142491a43d4..ad762b5ff56a6 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc @@ -263,6 +263,21 @@ variable should be passed via a dedicated argument (e.g. the value returned by `sub_group::get_local_id()[0]` could be passed as an integer to a `Callable` expecting a `sycl::ext::oneapi::experimental::simd`). +NOTE: Implementation must be able to calculate sub-group size based on the +`invoke_simd` SIMD call target and actual SPMD arguments. The following +algorithm should be employed: +- If any of the arguments is non-uniform, then the specification requires that +the corresponding formal SIMD argument must have `simd` or +`simd_mask` type. `N` is then the sub-group size. Implementation may +iterate through supported sub-group sizes at compile time and use +`std::is_invocable` to see if particular sub-group size fits. There must be +exactly one fitting sub-group size (matching `N`) or the user program is not +well-formed. +- Otherwise, if the return type of the SIMD target function is `simd` or +`simd_mask`, then the sub-group size is `N`. +- Otherwise (all arguments and the return type are uniform) the SIMD target is +fully uniform and it can be invoked with any sub-group size. + The `invoke_simd` function has the same requirements as other group functions (as defined in Section 4.17.3 of the SYCL 2020 specification). A call to `invoke_simd` must be encountered in converged control flow by all work-items From 9c6a72186df6058a5945ce67065cb78236581aac Mon Sep 17 00:00:00 2001 From: Konstantin S Bobrovsky Date: Mon, 15 Aug 2022 16:09:54 -0700 Subject: [PATCH 2/5] Add revision history line. --- .../extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc index ad762b5ff56a6..135ec65a014da 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc @@ -462,4 +462,5 @@ SYCL 2020 specification states that this only applies to inter-device transfers. |2|2021-03-31|John Pennycook|*Rename extension and add feature test macro* |3|2021-04-23|John Pennycook|*Split uniform wrapper into separate extension* |4|2022-01-20|John Pennycook|*Clarify interaction with SYCL_EXTERNAL* +|5|2022-08-15|Konst Bobrovskii|*Clarify sub-group size calculation* |======================================== From f741469073cea650f7b247376db49e0bfeabc02c Mon Sep 17 00:00:00 2001 From: kbobrovs Date: Fri, 19 Aug 2022 15:36:32 -0700 Subject: [PATCH 3/5] Update sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc Co-authored-by: rolandschulz --- .../sycl_ext_oneapi_invoke_simd.asciidoc | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc index 135ec65a014da..6fb560b5c4bbe 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc @@ -263,15 +263,15 @@ variable should be passed via a dedicated argument (e.g. the value returned by `sub_group::get_local_id()[0]` could be passed as an integer to a `Callable` expecting a `sycl::ext::oneapi::experimental::simd`). -NOTE: Implementation must be able to calculate sub-group size based on the -`invoke_simd` SIMD call target and actual SPMD arguments. The following -algorithm should be employed: -- If any of the arguments is non-uniform, then the specification requires that -the corresponding formal SIMD argument must have `simd` or -`simd_mask` type. `N` is then the sub-group size. Implementation may -iterate through supported sub-group sizes at compile time and use -`std::is_invocable` to see if particular sub-group size fits. There must be -exactly one fitting sub-group size (matching `N`) or the user program is not +NOTE: An implementation must compare the sub-group size to the `N` value of the SIMD +arguments of the SIMD function object to emit proper diagnostic. +The following algorithm may be employed: +- If any of the arguments is non-uniform, +the corresponding SIMD argument must have `simd` or +`simd_mask` type. `N` must be identical to the the sub-group size. Implementation may +iterate through the supported sub-group sizes at compile time and check whether the +function object is invocable with a sub-group size . There must be +exactly one sub-group size (matching `N`) which allows the function object to be invoked or the user program is not well-formed. - Otherwise, if the return type of the SIMD target function is `simd` or `simd_mask`, then the sub-group size is `N`. From c01393de64b572a1e81dfd89df921d7bf534abf2 Mon Sep 17 00:00:00 2001 From: Konstantin S Bobrovsky Date: Fri, 19 Aug 2022 17:00:57 -0700 Subject: [PATCH 4/5] Address review comments, clarify returning uniform values. --- .../sycl_ext_oneapi_invoke_simd.asciidoc | 39 +++++++++++++------ 1 file changed, 27 insertions(+), 12 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc index 6fb560b5c4bbe..0ce319a5537e8 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc @@ -123,10 +123,15 @@ of a `simd` type (e.g. using multiple registers rather than an array). We expect implementations of this ABI to build on Clang's `ext_vector_type`, rather than specifying the register types directly. -The value of `N` for each `simd` argument to a SIMD function must be the same, -and represents the sub-group size of the calling kernel. If a function accepts -`simd` arguments with multiple `N` values, it cannot be invoked from a SPMD -kernel and the implementation must issue a diagnostic in this case. +Each of the following must hold true for a SIMD function and its invoke_simd +call site: +- the value of `N` of each `simd` formal parameter corresponding to a non-uniform +SPMD actual argument (if any) must be the same +- if the return type of the SIMD function is `simd` or `simd_mask`, + then `M` must match `N` above +Otherwise, the SIMD function cannot be invoked from an SPMD kernel and the +implementation must issue a diagnostic in this case. +The `N` value represents the sub-group size of the calling kernel. The example below shows a simple SIMD function that scales all elements of an 8-wide SIMD type `x` by a scalar value `n`. @@ -244,6 +249,12 @@ rules in reverse: are converted to `T`, and the value in element `i` of the SIMD type is returned to the work-item with sub-group local ID `i`. +- Return values of type `sycl::ext::oneapi::experimental::uniform` are not anyhow converted, + and broadcast to each work-item; every work-item in the sub-group receives + the same value. + NOTE: `sycl::ext::oneapi::experimental::uniform` return type is the way to return + a uniform value of `simd` or `simd_mask` type. + - Return values of type `T` are converted to `sycl::ext::oneapi::experimental::uniform`, and broadcast to each work-item; every work-item in the sub-group receives the same value. @@ -263,15 +274,19 @@ variable should be passed via a dedicated argument (e.g. the value returned by `sub_group::get_local_id()[0]` could be passed as an integer to a `Callable` expecting a `sycl::ext::oneapi::experimental::simd`). -NOTE: An implementation must compare the sub-group size to the `N` value of the SIMD -arguments of the SIMD function object to emit proper diagnostic. -The following algorithm may be employed: -- If any of the arguments is non-uniform, -the corresponding SIMD argument must have `simd` or -`simd_mask` type. `N` must be identical to the the sub-group size. Implementation may +NOTE: An implementation must be able to determine the sub-group size of the SIMD +function object to perform argument type conversion. +The determined subgroup size must be equal to the value of the caller's `sycl::reqd_sub_group_size` +attribute (if present) and consistent with the argument conversion rules defined +by this specification. Otherwise, the implementation must issue a proper diagnostic. + +The following algorithm to determine the sub-group size may be employed: +- If any of the actual SPMD arguments is non-uniform, the specification requires +that the corresponding SIMD argument must have `simd` or `simd_mask` +type, and `N` must be equal to the the sub-group size. The implementation may iterate through the supported sub-group sizes at compile time and check whether the -function object is invocable with a sub-group size . There must be -exactly one sub-group size (matching `N`) which allows the function object to be invoked or the user program is not +function object is invocable with given sub-group size. There must be +exactly one such sub-group size and it must match `N` or the user program is not well-formed. - Otherwise, if the return type of the SIMD target function is `simd` or `simd_mask`, then the sub-group size is `N`. From 06f7d360905dd406bdefe35e248cdf4d558eee56 Mon Sep 17 00:00:00 2001 From: kbobrovs Date: Tue, 23 Aug 2022 15:56:39 -0700 Subject: [PATCH 5/5] Update sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc --- .../experimental/sycl_ext_oneapi_invoke_simd.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc index 0ce319a5537e8..d6676108a67b2 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc @@ -125,7 +125,7 @@ rather than specifying the register types directly. Each of the following must hold true for a SIMD function and its invoke_simd call site: -- the value of `N` of each `simd` formal parameter corresponding to a non-uniform +- the value of `N` of each `simd` or `simd_mask` formal parameter corresponding to a non-uniform SPMD actual argument (if any) must be the same - if the return type of the SIMD function is `simd` or `simd_mask`, then `M` must match `N` above