CP016: Sub groups proposal #75
Conversation
* Add initial proposal for adding sub groups to SYCL.
* Add links to OpenCL 2.2 and OpenCL C++ 1.0 specifications. * Add specification of the codeplay::nd_item class template. * Add proposal to readme. * Add minor formatting changes.
* Move definition of codeplay extension nd_item to a class definition. * Add namespaces for class definitions.
…/standards-proposals into sub-groups-proposal
…ds-proposals into sub-groups-proposal
…/standards-proposals into sub-groups-proposal
* Add description of sub groups example.
…ds-proposals into sub-groups-proposal
…/standards-proposals into sub-groups-proposal
* Removed unnecessary namespace
| ## Overview | ||
|
|
||
| This vendor extension aims to define an interface to expose sub-group functionality, | ||
| as defined in the SYCL 2.2 provisional and the OpenCL 2.2 provisional, |
There was a problem hiding this comment.
OpenCL 2.2 isn't provisional (could probably also say "OpenCL 2.1" here as that is when subgroups were added to the main spec).
| @@ -0,0 +1,157 @@ | |||
| # Basic Sub group support | |||
There was a problem hiding this comment.
OpenCL uses "Sub-group", not "Sub group".
This is on many lines, I won't comment on them all.
|
|
||
| All new functionality is exposed under the `basic_sub_group` namespace | ||
| in the `codeplay` vendor extension namespace. | ||
| When the vendor extension `basic_sub_group` is available, the macro |
There was a problem hiding this comment.
OpenCL subgroup extensions used subgroups in their names (e.g. cl_khr_subgroups).
It may be valid to ignore that given the type cl::sycl::codeplay::sub_group.
| sub-groups, with some restrictions: | ||
|
|
||
| * The number of sub-groups available for each work-group is determined | ||
| at compile-time and remains the same during the execution of the SYCL application. |
There was a problem hiding this comment.
I'm not as familiar with the details of the SYCL spec, but looking at it now it seems that current use of "compile-time" means when compiling the SYCL program rather than the SYCL runtime calling clCompileProgram or similar. This information can't be known until clGetKernelSubGroupInfo (or clGetKernelSubGroupInfoCODEPLAY) can be called.
I'm just unclear on the intended meaning of compile-time.
|
|
||
| * The number of sub-groups available for each work-group is determined | ||
| at compile-time and remains the same during the execution of the SYCL application. | ||
| * The number of threads per sub-group is known at compile-time, and remains the |
There was a problem hiding this comment.
This use of "threads" seems unusual.
"The sub-group size is known ...`
perhaps?
| namespace sycl { | ||
| namespace codeplay { | ||
|
|
||
| template <int Dimensions> |
There was a problem hiding this comment.
I see that this comes from the SYCL 2.2 provisional specification, but for my own sanity: Dimensions is here as hypothetical future proofing and barring some future vendor extension it will always have the value 1?
There was a problem hiding this comment.
Or is it important that, for example, with a 3D ND-Range that sub-group also has Dimension of 3, even if the actual sizes are always of the form {N, 1, 1}?
| @@ -0,0 +1,36 @@ | |||
| # Basic sub-group extension | |||
There was a problem hiding this comment.
General comment: Seems like this proposal is missing device info properties?
info::device::max_num_sub_groups and info::device::sub_group_independent_forward_progress
|
|
||
| ## Namespace `basic_sub_group` | ||
|
|
||
| All new functionality is exposed under the `basic_sub_group` namespace |
There was a problem hiding this comment.
There doesn't seem to be any basic_sub_group namespace used in this extension?
| sub-groups, with some restrictions: | ||
|
|
||
| * The number of sub-groups available for each work-group is determined | ||
| at compile-time and remains the same during the execution of the SYCL application. |
There was a problem hiding this comment.
compile-time is this unclear: SYCL compile time or OpenCL.
The text after seems to imply SYCL because it sub-group range is constexpr, but this could be made more explicit here.
| class sub_group { | ||
| public: | ||
|
|
||
| constexpr range<Dimensions> get_sub_group_range() const; |
There was a problem hiding this comment.
I find this really restrictive, you basically ban compilation for generic targets. But I also understand the need if you know your underlying target. Could we have a in-between solution ? i.e. constexpr iff you know the underlying target and it properties
Add a proposal for introducing an extension to support sub-groups in SYCL 1.2.1.