From adfdc4ccdc8664b36bf8b56dc3788ac8b5874f74 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Peter=20=C5=BDu=C5=BEek?= Date: Sat, 21 Sep 2019 09:53:57 +0100 Subject: [PATCH 1/3] Accessor aliases In order to reduce the verbosity of programming with SYCL accessors, this proposal aims to reduce the number of template parameters to just 2 from the current 4 (placeholders accessors even have 5). It achieves this by adding aliases to accessors based on the access target and slightly revised rules for access modes. 1. Added aliases to accessors * `buffer_accessor`, * `constant_buffer_accessor * `host_accessor` 2. Deduce access mode based on constness of data type * `access::mode::read` for `const dataT` * `access::mode::read_write` for `dataT` 3. Allow conversions to new aliases 4. Return aliases from buffer access * `get_device_access` * `get_device_constant_access` * `get_host_access` 5. Allow any accessor to be registered with `handler::require` 6. Overload for `handler::require` that also takes an access mode * Hint to the scheduler 7. Type traits to help with deducing access modes * `access_mode_from_type` * `type_from_access_mode` 8. Default all accessor template parameters * Expect the data type * Default to an accessor to global buffer with read-write access * Assumes any accessor can be a placeholder 9. Discussed some considerations and alternatives 10. Examples of reduced verbosity --- README.md | 3 +- accessor-alias/index.md | 20 ++ accessor-alias/sycl-2.2/index.md | 571 +++++++++++++++++++++++++++++++ 3 files changed, 593 insertions(+), 1 deletion(-) create mode 100644 accessor-alias/index.md create mode 100644 accessor-alias/sycl-2.2/index.md diff --git a/README.md b/README.md index 5099833..3171b8a 100644 --- a/README.md +++ b/README.md @@ -63,4 +63,5 @@ from this registry in the future. | CP020 | [Interop Task](interop_task/interop_task.md) | SYCL 1.2.1 | 16 January 2019 | 16 January 2019 | _Available since CE 1.0.5_ | | CP021 | [Default-Constructed Buffers](default-constructed-buffers/default-constructed-buffers.md) | SYCL 1.2.1 | 27 August 2019 | 5 September 2019 | _Draft_ | | CP022 | [Host Task with Interop capabilities](host_task/host_task.md) | SYCL 1.2.1 | 16 January 2019 | 20 January 2020 | _Final Draft_ | -| CP026 | [Generalized Error Handling For SYCL](error-handling/sycl-error-handling.md) | SYCL Next | 10 March 2020 | 10 March 2020 | _Under Review_ | \ No newline at end of file +| CP023 | [Accessor aliases](accessor-alias/index.md) | SYCL 1.2.1 vendor extension | 22 September 2019 | 5 March 2020 | _Work in Progress_ | +| CP026 | [Generalized Error Handling For SYCL](error-handling/sycl-error-handling.md) | SYCL Next | 10 March 2020 | 10 March 2020 | _Under Review_ | diff --git a/accessor-alias/index.md b/accessor-alias/index.md new file mode 100644 index 0000000..8a6615f --- /dev/null +++ b/accessor-alias/index.md @@ -0,0 +1,20 @@ +# Accessor aliases + +| Proposal ID | CP023 | +|-------------|--------| +| Name | Accessor aliases | +| Date of Creation | 22 September 2019 | +| Target | SYCL 1.2.1 vendor extension | +| Current Status | _Work in Progress_ | +| Reply-to | Peter Žužek | +| Original author | Peter Žužek | +| Contributors | | + +## Overview + +This paper proposes the addition of target specific aliases to accessors +in order to reduce accessor verbosity. + +## Versions + +[Version 1](sycl-2.2/index.md) diff --git a/accessor-alias/sycl-2.2/index.md b/accessor-alias/sycl-2.2/index.md new file mode 100644 index 0000000..c09bc13 --- /dev/null +++ b/accessor-alias/sycl-2.2/index.md @@ -0,0 +1,571 @@ +# Accessor aliases + +## Motivation + +Accessors are the cornerstone of SYCL +as they are used inside and outside kernels to access data +in a pointer-like manner. +However, SYCL ecosystem projects have shown that +they can be a bit awkward to handle due to their verbosity. +Storing accessors in functors requires specifying all template parameters +and accessors of different access modes are treated as different types, +even though there isn't much difference in accessing data. + +### Existing `local_accessor` alias + +SYCL 1.2.1 already has an alias for local memory, +which can serve as an illustrative example of an alias reducing verbosity. + +```cpp +// Without the alias +accessor + localAccOld{range<1>{32}}; + +// With the alias +local_accessor + localAcc{range<1>{32}}; +``` + +## Summary + +This proposal introduces access-target-specific aliases +for accessors. +In order to simplify access modes, +the aliases use only one of two modes: +`access::mode::read` for `const dataT` data parameters +and `access::mode::read_write` for non-const `dataT` data parameters. +This allows each alias to have only two template parameters, +`dataT` and `dims`, where `dims` is also defaulted to `1`. + +The new aliases are: +* `buffer_accessor` +* `constant_buffer_accessor` +* `host_accessor` + +We also propose defaulting all accessor template parameters, +except for the type parameter. + +The proposal slightly changes the semantics of accessing data, +but still aims to be completely backwards compatible with SYCL 1.2.1. + +Note that this proposal assumes that any accessor can be a placeholder, +essentially deprecating the `access::target::placeholder` template parameter. +This is covered by a separate proposal. +It also doesn't address image accessors. +We discuss this and some other issues +in the [Considerations and alternatives](#considerations-and-alternatives) section. + +## Simplifying access modes + +We can think of an accessor as performing two basic functions: +1. Requesting data access +2. Providing access to data + +When requesting access to a buffer, +the user has to specify the access mode. +There are 6 access modes in the SYCL 1.2.1 specification, +and selecting the right modes when requesting access +can provide a lot of information to the scheduler. + +However, when it comes to the second function of the accessor, +providing access to data, +the plethora of access modes is less useful. +If we consider a "standard", non-fancy C++ pointer, +it can either provide read and write access to the underlying data, +or read-only access (pointer-to-const). +It cannot provide write-only access - +that kind of functionality would be reserved to only niche applications anyway. + +We would like accessors to behave in a similar manner. +As mentioned, access modes are still very useful to the scheduler, +so we would not deprecate any of them. +Instead, we propose that when it comes to providing data access, +there are only two relevant access modes: +1. `access::mode::read` +2. `access::mode::read_write` + +This means that accessors would only provide different APIs +(in terms of accessing data) +for these two access modes, +and any other mode would resolve to one of these two. + +Additionally, we propose that these two access modes +are directly tied to the constness of the data type. +So `access::mode::read` would be tied to `const dataT` +and "discourage" a non-const `dataT`, +while `access::mode::read_write` would be tied to a non-const `dataT` +and "discourage" `const dataT`. +The "discouraging" part isn't quite obvious +and we discuss it a bit more in the +[Discouraging access mode combinations](#discouraging-access-mode-combinations) +section. + + +### Type traits for resolving access modes + +In order to help with resolving access modes and the constness of types, +we propose adding a few type traits: + +```cpp +namespace cl { +namespace sycl { + +/// access_mode_constant +template +using access_mode_constant = + std::integral_constant; + +/// access_mode_from_type +template +struct access_mode_from_type + : access_mode_constant {}; +template +struct access_mode_from_type + : access_mode_constant {}; + +/// type_from_access_mode +template +struct type_from_access_mode { + using type = dataT; + static_assert( + (!std::is_const::value || (requestedMode == access::mode::read)), + "Cannot request write access to const data"); +}; +template +struct type_from_access_mode { + using type = const dataT; +}; +template +using type_from_access_mode_t = + typename type_from_access_mode::type; + +} // namespace sycl +} // namespace cl +``` + +| Type trait | Description | +|-----------------|-------------| +| *`template access_mode_constant`* | Alias that stores `access::mode` into `std::integral_constant`. | +| *`template access_mode_from_type`* | Deduces access mode based on the constness of `dataT`. | +| *`template type_from_access_mode`* | Deduces the constness of `dataT` based on the access mode. Fails when requesting write access on `const dataT`. | + +## Defaulting template parameters + +We propose defaulting all accessor template parameters +except for the type parameter. +Defaulting to a read-only 1-dimensional global buffer accessor +reduces a large amount of verbosity for the simplest cases +and makes it easier to prototype SYCL code. + +```cpp +namespace cl { +namespace sycl { + +template < + typename dataT, + int dims = 1, + access::mode accMode = access::mode::read_write, + access::target accTarget = access::target::global_buffer, + access::placeholder isPlaceholder = access::placeholder::false_t> +class accessor; + +} // namespace sycl +} // namespace cl +``` + +## Aliases + +We propose the following aliases to accessors: + +```cpp +namespace cl { +namespace sycl { + +template +using buffer_accessor = + accessor::value, + access::target::global_buffer>; + +template +using constant_buffer_accessor = + accessor; + +template +using host_accessor = + accessor::value, + access::target::host_buffer>; + +} // namespace sycl +} // namespace cl +``` + +## Explicit conversions to aliases + +We propose allowing an accessor +that isn't quite the same type as one of the aliases - +for example, by having a `write` instead of a `read_write` access mode - +to be explicitly converted to an alias type. + +```cpp +namespace cl { +namespace sycl { + +template < + typename dataT, + int dims = 1, + access::mode accMode = access::mode::read_write, + access::target accTarget = access::target::global_buffer, + access::placeholder isPlaceholder = access::placeholder::false_t> +class accessor { + public: + /// All existing members here + + ... + + // Explicit conversion to `buffer_accessor` + // Only allowed when `accTarget == access::target::global_buffer` + explicit operator + buffer_accessor< + type_from_access_mode_t, dims>(); + + // Explicit conversion to `constant_buffer_accessor` + // Only allowed when `accTarget == access::target::constant_buffer` + explicit operator + constant_buffer_accessor< + type_from_access_mode_t, dims>(); + + // Explicit conversion to `host_accessor` + // Only allowed when `accTarget == access::target::host_buffer` + explicit operator + host_accessor< + type_from_access_mode_t, dims>(); +}; + +} // namespace sycl +} // namespace cl +``` + +| Member function | Description | +|-----------------|-------------| +| *`explicit operator buffer_accessor, dims>()`* | Performs a cast to `buffer_accessor`. Only allowed when `accTarget == access::target::global_buffer` and if `accMode` doesn't discard the constness of `dataT`. | +| *`explicit operator constant_buffer_accessor()`* | Performs a cast to `constant_buffer_accessor`. Only allowed when `accTarget == access::target::constant_buffer`. | +| *`explicit operator host_accessor, dims>()`* | Performs a cast to `host_accessor`. Only allowed when `accTarget == access::target::host_buffer` and if `accMode` doesn't discard the constness of `dataT`. | + +## Extending buffer class to return aliases + +In order to maintain backwards compatibility with SYCL 1.2.1, +the `get_access` member function should not change, +apart from defaulting the template parameters. +Instead, we propose adding new member functions to the `buffer` class +that request data access and return one of the new aliases. + +```cpp +namespace cl { +namespace sycl { + +template +class buffer { + public: + /// All existing members here + + ... + + + /// Existing `get_access` gains defaulted parameters + + template + accessor + get_access(handler& cgh); + + template + accessor + get_access(); + + + /// New functions + + template + buffer_accessor< + type_from_access_mode_t, dims> + get_device_access(handler& cgh); + + constant_buffer_accessor + get_device_constant_access(handler& cgh); + + template + host_accessor< + type_from_access_mode_t, dims> + get_host_access(); +}; + +} // namespace sycl +} // namespace cl +``` + +| Member function | Description | +|-----------------|-------------| +| *`template buffer_accessor, dims> get_device_access(handler& cgh)`* | Calls `get_access`, but returns a `buffer_accessor`. | +| *`constant_buffer_accessor get_device_constant_access(handler& cgh)`* | Calls `get_access`, but returns a `constant_buffer_accessor`. | +| *`template host_accessor, dims> get_host_access()`* | Calls `get_access`, but returns a `host_accessor`. | + +## Extending the handler + +Reducing effective access modes from 6 to 2 +would have some impact on the scheduler +which could now have less information to guide the scheduling process. +This would be especially problematic for placeholder accessors, +where the type of the accessor is determined +before they are registered with a command group. +To resolve this, we propose extending the `handler` class +by allowing `require` to be called on any accessor +and to add an overload of `require` that also takes an access mode. + +An example would be to call `require` with `global_accessor` +and `access::mode::discard_read_write`, +which would inform the scheduler to not copy over any old data. + +```cpp +namespace cl { +namespace sycl { + +class handler { + public: + /// Existing functions + /// ... + + /// New functions + + // Registers an accessor with a command group submission + template + void require(accessor acc); + + // Registers an accessor with a command group submission + // `requestedMode` is a hint to the scheduler + template + void require(accessor acc, + access::mode requestedMode); +}; + +} // namespace sycl +} // namespace cl +``` + +| Member function | Description | +|-----------------|-------------| +| *`template void require(accessor acc)`* | Registers the accessor for command group submission. | +| *`template void require(accessor acc, access::mode requestedMode)`* | Registers the accessor for command group submission. `requestedMode` is used as a hint to the scheduler. `requestedMode` cannot be a write mode if the accessor mode is `access::mode::read`. | + +## Considerations and alternatives + +### Image accessors + +Images are different from buffers +in the sense that they have three basic access modes +instead of just two for buffer: +in addition to being read-only and read-write, +they can also be write-only. +Additionally, read-write images are not supported on all devices, +this is even an extension in OpenCL. +This prevents us from using the constness of the data type, +which only has two states, +to determine the access mode. +Any kind of image accessor alias +would thus need to incorporate the access mode as a template parameter, +reducing the usability of such alias. + +Image array access is another area of contention +and we are not sure what approach would best suit +in reducing verbosity there. + +### Placeholder accessors + +This proposal builds on top of allowing any accessor to be a placeholder. +The proposal still works without that assumption, +but its usefulness is significantly reduced. +If any accessor can be a placeholder, +then the `isPlaceholder` template parameter becomes obsolete, +making it easy to reduce the 5 template parameters of the `accessor` +to just 2 for each alias. +Without that change, +placeholder accessors would not benefit from this proposal, +unless some changes are made. + +Here are some options: +1. Add an `isPlaceholder` parameter to the aliases, default it to `false_t`. + This wouldn't reduce accessor verbosity as much as originally planned, + by still requiring three template parameters for placeholder accessor aliases. +2. Introduce more aliases based on whether the accessor is a placeholder or not. + We don't consider this option very desirable + since it just replaces one kind of verbosity for another. + +### Discouraging access mode combinations + +It is not clear to what extent should the specification discourage +the use of a non-const accessor `dataT` with `access::mode::read` +and a const accessor `dataT` with a write mode. +We considered deprecating those use cases, +or maybe just adding a warning, +or even `static_assert` to prevent their usage. +However, at least the example of `access::mode::read` with a non-const `dataT` +is very common in existing SYCL code +because using `const` is not considered essential in SYCL 1.2.1. + +## Examples + +### Storing accessors in a kernel functor + +There is an example in the +[ComputeCpp SDK](https://github.com/codeplaysoftware/computecpp-sdk/blob/master/samples/template-function-object.cpp) +that showcases a kernel functor with stored accessors. +It uses user-defined aliases `read_accessor` and `write_accessor` +to simplify the three private members - +two read accessors and one write accessor. +These kinds of user-defined aliases seem to be a regular occurrence +in SYCL ecosystem code. +As shown by the example below, +this proposal would essentially standardize the different aliases +from across different projects. + +```cpp +using namespace cl::sycl; + +template +class vector_add_kernel { + public: + vector_add_kernel(buffer_accessor ptrA, + buffer_accessor ptrB, + buffer_accessor ptrC) + : m_ptrA(ptrA), m_ptrB(ptrB), m_ptrC(ptrC) {} + + void operator()(item<1> item) { + m_ptrC[item.get_id()] = m_ptrA[item.get_id()] + m_ptrB[item.get_id()]; + } + + private: + buffer_accessor m_ptrA; + buffer_accessor m_ptrB; + buffer_accessor m_ptrC; +}; + +``` + +### New `get_access` functions + +```cpp +using namespace cl::sycl; + +// Assume these are available +handler cgh; +buffer buf; + +// Global buffer access +accessor bufAcc = + buf.get_access(cgh); +buffer_accessor bufAccNew = + buf.get_device_access(cgh); + +// Global buffer access, read-only +accessor + bufAccRead = buf.get_access(cgh); +buffer_accessor bufAccReadNew = + buf.get_device_access(cgh); + +// Global buffer access, ignore previous data +accessor + bufAccDiscard = buf.get_access(cgh); +buffer_accessor bufAccNewDiscard = + buf.get_device_access(cgh); + +// Constant buffer access +accessor + bufAccConst = buf.get_access(cgh); +constant_buffer_accessor bufAccConstNew = + buf.get_device_constant_access(cgh); + +// Host buffer access +accessor + bufAccHost = buf.get_access(); +host_accessor bufAccHostNew = + buf.get_host_access(); + +// Host buffer access, read-only +accessor + bufAccHostRead = buf.get_access(); +host_accessor bufAccHostNewRead = + buf.get_host_access(); + +// Host buffer access, ignore previous data +accessor + bufAccHostDiscard = buf.get_access(); +host_accessor bufAccHostNewDiscard = + buf.get_host_access(); +``` + +### Calling `require` + +```cpp +using namespace cl::sycl; + +// Assume these are available +queue q; +buffer buf1; +buffer buf2; + +q.submit([](handler& cgh) { + // Request read-write access to both buffers + buffer_accessor bufAcc1 = + buf1.get_device_access(cgh); + buffer_accessor bufAcc2 = + buf2.get_device_access(cgh); + + // Register the accessor for command group submission + // Not so useful in this case since it's already been registered, + // but bufAcc1 could also be a placeholder + cgh.require(bufAcc1); + + // Register the accessor for command group submission + // bufAcc2 has also already been registered, + // but this instructs the scheduler to ignore previous data + cgh.require(bufAcc2, access::mode::discard_read_write); + + ... +}); + +``` From 736ef71c10e9a1b1df14a3b242d466d35185b582 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Peter=20=C5=BDu=C5=BEek?= Date: Thu, 5 Mar 2020 10:49:12 +0000 Subject: [PATCH 2/3] CP023 Accessor aliases, revision 0.2 In order to reduce the verbosity of programming with SYCL accessors, this proposal aims to reduce the number of template parameters to just 2 from the current 5. It achieves this by slightly revising how to treat read-only data and adding alias templates based on the access target. 1. Main changes * Default accessor template parameters * Simplify access modes * Alias templates based on access target * Extend the handler class 2. Accessor alias templates * `constant_buffer_accessor * `host_accessor` 3. Treat `const T` the same as `access::mode::read` * Simplifies a lot of code 4. Define implicit conversions for equivalent types 5. Implicit conversions that add `const` 6. Overload for `handler::require` that also takes an access mode * To weaken the access mode 8. Default all accessor template parameters * Expect the data type * Default to an accessor to global buffer with read-write access * Assumes accessors can be a placeholder without template parameter 9. Discussed some considerations and alternatives 10. Examples of reduced verbosity --- README.md | 2 +- accessor-alias/index.md | 6 +- accessor-alias/sycl-2.2/index.md | 706 ++++++++++++++++++------------- 3 files changed, 419 insertions(+), 295 deletions(-) diff --git a/README.md b/README.md index 3171b8a..aa3dfc0 100644 --- a/README.md +++ b/README.md @@ -63,5 +63,5 @@ from this registry in the future. | CP020 | [Interop Task](interop_task/interop_task.md) | SYCL 1.2.1 | 16 January 2019 | 16 January 2019 | _Available since CE 1.0.5_ | | CP021 | [Default-Constructed Buffers](default-constructed-buffers/default-constructed-buffers.md) | SYCL 1.2.1 | 27 August 2019 | 5 September 2019 | _Draft_ | | CP022 | [Host Task with Interop capabilities](host_task/host_task.md) | SYCL 1.2.1 | 16 January 2019 | 20 January 2020 | _Final Draft_ | -| CP023 | [Accessor aliases](accessor-alias/index.md) | SYCL 1.2.1 vendor extension | 22 September 2019 | 5 March 2020 | _Work in Progress_ | +| CP023 | [Accessor aliases](accessor-alias/index.md) | SYCL Next (after 1.2.1) | 22 September 2019 | 31 March 2020 | _Work in Progress_ | | CP026 | [Generalized Error Handling For SYCL](error-handling/sycl-error-handling.md) | SYCL Next | 10 March 2020 | 10 March 2020 | _Under Review_ | diff --git a/accessor-alias/index.md b/accessor-alias/index.md index 8a6615f..8811e5e 100644 --- a/accessor-alias/index.md +++ b/accessor-alias/index.md @@ -4,7 +4,9 @@ |-------------|--------| | Name | Accessor aliases | | Date of Creation | 22 September 2019 | -| Target | SYCL 1.2.1 vendor extension | +| Revision | 0.2 | +| Latest Update | 31 March 2020 | +| Target | SYCL Next (after 1.2.1) | | Current Status | _Work in Progress_ | | Reply-to | Peter Žužek | | Original author | Peter Žužek | @@ -17,4 +19,4 @@ in order to reduce accessor verbosity. ## Versions -[Version 1](sycl-2.2/index.md) +[Revision 0.2](sycl-2.2/index.md) diff --git a/accessor-alias/sycl-2.2/index.md b/accessor-alias/sycl-2.2/index.md index c09bc13..fefbb08 100644 --- a/accessor-alias/sycl-2.2/index.md +++ b/accessor-alias/sycl-2.2/index.md @@ -1,5 +1,19 @@ # Accessor aliases +## Table of Contents + +* [Motivation](#motivation) +* [Revisions](#revisions) +* [Summary](#summary) +* Changes + * [Defaulting template parameters](#defaulting-template-parameters) + * [Simplifying access modes](#simplifying-access-modes) + * [Aliases](#aliases) + * [Extending the handler](#extending-the-handler) +* [Considerations and alternatives](#considerations-and-alternatives) + +### TLDR: [Examples](#examples) + ## Motivation Accessors are the cornerstone of SYCL @@ -7,7 +21,7 @@ as they are used inside and outside kernels to access data in a pointer-like manner. However, SYCL ecosystem projects have shown that they can be a bit awkward to handle due to their verbosity. -Storing accessors in functors requires specifying all template parameters +Storing accessors in function objects requires specifying all template parameters, and accessors of different access modes are treated as different types, even though there isn't much difference in accessing data. @@ -29,39 +43,81 @@ local_accessor localAcc{range<1>{32}}; ``` +## Revisions + +### 0.2 + +* Simplified handling of non-writeable accessors +* Defined implicit conversions w.r.t. constness and access modes +* Removed `buffer_accessor` alias +* Removed type traits +* Removed member functions for returning aliases +* Added `access::mode` template parameter to `host_accessor` +* Removed section _Discouraging access mode combinations_ +* Discuss CTAD +* Discuss trait for writeable accessors +* More and better examples +* Added table of contents +* Deprecate `access::mode::atomic` as part of this proposal +* Link to proposal that deprecates `accessor::placeholder` +* Minor fixes + +### 0.1 + +* Initial proposal + ## Summary -This proposal introduces access-target-specific aliases -for accessors. -In order to simplify access modes, -the aliases use only one of two modes: -`access::mode::read` for `const dataT` data parameters -and `access::mode::read_write` for non-const `dataT` data parameters. -This allows each alias to have only two template parameters, -`dataT` and `dims`, where `dims` is also defaulted to `1`. - -The new aliases are: -* `buffer_accessor` -* `constant_buffer_accessor` -* `host_accessor` - -We also propose defaulting all accessor template parameters, -except for the type parameter. +Main points of the proposal: + +1. [Default `accessor` template parameters](#defaulting-template-parameters) +1. [Clarify and simplify access modes](#simplifying-access-modes) + * `read` and `read_write` most important + * Allow `const T` to denote read-only data access + * Deprecate `access::mode::atomic` +1. [Introduce `access::target`-specific aliases for accessors](#aliases) +1. [Extending `handler::require`](#extending-the-handler) The proposal slightly changes the semantics of accessing data, but still aims to be completely backwards compatible with SYCL 1.2.1. Note that this proposal assumes that any accessor can be a placeholder, -essentially deprecating the `access::target::placeholder` template parameter. -This is covered by a separate proposal. +essentially deprecating the `access::placeholder` template parameter. +This is covered by a separate proposal: +[CP024 Default placeholders](https://github.com/codeplaysoftware/standards-proposals/pull/122). It also doesn't address image accessors. We discuss this and some other issues in the [Considerations and alternatives](#considerations-and-alternatives) section. +## Defaulting template parameters + +We propose defaulting all accessor template parameters +except for the type parameter. +Defaulting to a read-only 1-dimensional global buffer accessor +reduces a large amount of verbosity for the simplest cases +and makes it easier to prototype SYCL code. + +```cpp +namespace cl { +namespace sycl { + +template < + typename dataT, + int dims = 1, + access::mode accMode = access::mode::read_write, + access::target accTarget = access::target::global_buffer, + access::placeholder isPlaceholder = access::placeholder::false_t> +class accessor; + +} // namespace sycl +} // namespace cl +``` + ## Simplifying access modes We can think of an accessor as performing two basic functions: -1. Requesting data access + +1. Requesting access to data 2. Providing access to data When requesting access to a buffer, @@ -84,6 +140,7 @@ As mentioned, access modes are still very useful to the scheduler, so we would not deprecate any of them. Instead, we propose that when it comes to providing data access, there are only two relevant access modes: + 1. `access::mode::read` 2. `access::mode::read_write` @@ -92,105 +149,186 @@ This means that accessors would only provide different APIs for these two access modes, and any other mode would resolve to one of these two. -Additionally, we propose that these two access modes -are directly tied to the constness of the data type. -So `access::mode::read` would be tied to `const dataT` -and "discourage" a non-const `dataT`, -while `access::mode::read_write` would be tied to a non-const `dataT` -and "discourage" `const dataT`. -The "discouraging" part isn't quite obvious -and we discuss it a bit more in the -[Discouraging access mode combinations](#discouraging-access-mode-combinations) -section. +Additionally, we propose that the constness of the data type be considered +when providing data access. +As a final step for simplifying access modes, +we propose `access::mode::atomic` be deprecated. +It is already possible to construct a `cl::sycl::atomic` +by using the pointer obtained from an accessor. -### Type traits for resolving access modes +### Constness of data -In order to help with resolving access modes and the constness of types, -we propose adding a few type traits: +Existing SYCL 1.2.1 code usually uses +non-`const`-qualified data types for accessor. +(For examples please take a look at SYCL ecosystem projects, +like SYCL-BLAS, SYCL-DNN, or the SYCL backend of TensorFlow.) +It might define it's own aliases to read and write accessors +when trying to store accessors in a kernel function object. -```cpp -namespace cl { -namespace sycl { +There is an example in the +[ComputeCpp SDK](https://github.com/codeplaysoftware/computecpp-sdk/blob/master/samples/template-function-object.cpp#L51) +that showcases a kernel functor with stored accessors. +It uses user-defined aliases `read_accessor` and `write_accessor` +to simplify the three private members - +two read accessors and one write accessor. +These kinds of user-defined aliases seem to be a regular occurrence +in SYCL ecosystem code. -/// access_mode_constant -template -using access_mode_constant = - std::integral_constant; +We propose using the constness of the data type +to simplify some of these use cases. +An example on how the above-linked example could be rewritten: -/// access_mode_from_type -template -struct access_mode_from_type - : access_mode_constant {}; -template -struct access_mode_from_type - : access_mode_constant {}; - -/// type_from_access_mode -template -struct type_from_access_mode { - using type = dataT; - static_assert( - (!std::is_const::value || (requestedMode == access::mode::read)), - "Cannot request write access to const data"); -}; +```cpp template -struct type_from_access_mode { - using type = const dataT; -}; -template -using type_from_access_mode_t = - typename type_from_access_mode::type; - -} // namespace sycl -} // namespace cl -``` +class vector_add_kernel { + public: + vector_add_kernel(accessor ptrA, + accessor ptrB, + accessor ptrC) + : m_ptrA(ptrA), m_ptrB(ptrB), m_ptrC(ptrC) {} -| Type trait | Description | -|-----------------|-------------| -| *`template access_mode_constant`* | Alias that stores `access::mode` into `std::integral_constant`. | -| *`template access_mode_from_type`* | Deduces access mode based on the constness of `dataT`. | -| *`template type_from_access_mode`* | Deduces the constness of `dataT` based on the access mode. Fails when requesting write access on `const dataT`. | + void operator()(item<1> item) { ... } -## Defaulting template parameters + private: + accessor m_ptrA; + accessor m_ptrB; + accessor m_ptrC; +}; +``` -We propose defaulting all accessor template parameters -except for the type parameter. -Defaulting to a read-only 1-dimensional global buffer accessor -reduces a large amount of verbosity for the simplest cases -and makes it easier to prototype SYCL code. +To enable this, we propose allowing `dataT` accessor template parameter +to be `const`, +but only when the access mode is `read` or `read_write`. +It's natural to allow it for the `read` access mode +since `const dataT` and `access::mode::read` +express essentially the same concept +when it comes to accessing data. + +The reason to allow `const dataT` for the `read_write` access mode +is because this is the default access mode, +which enables writing `accessor` in the above example. +The SYCL scheduler treats this combination +as if the access mode was `access::mode::read`. + +### Subscript operator + +After _requesting_ access to data by constructing an accessor, +`accessor::operator[]` is the function that enables +the second main functionality of the accessor class: +it _provides_ access to data. + +We propose that `accessor::operator[]` returns `reference_t`, +which is defined as: + +* `const T&` when the access mode is `access::mode::read` + or when `dataT` is `const`-qualified +* `T&` otherwise + +### Implicit conversions + +In order to simplify user code, +we propose allowing certain implicit conversions +that don't modify scheduling information. + +In standard C++ code, adding `const` qualifiers is almost always allowed. +In SYCL, going from a `read_write` accessor to a `read` accessor +is analogous to adding the `const` qualifier. +This allows us to have the following rules regarding `const`: + +1. Convert an accessor of `dataT` data type + to an accessor of `const dataT` data type, + all other template parameters being equal +1. Convert an accessor of `const dataT` data type + and `access::mode::read_write` mode + to an accessor of `dataT` and `access::mode::read` mode, + all other template parameters being equal +1. Convert an accessor of `const dataT` data type + and `access::mode::read_write` mode + to an accessor of `const dataT` and `access::mode::read` mode, + all other template parameters being equal +1. A combination of the previous rules, + allow any of these accessor combinations to be implicitly convertible + between each other: + * {`dataT`, `read`} + * {`const dataT`, `read`} + * {`const dataT`, `read_write`} + +An example of code these rules enable: ```cpp -namespace cl { -namespace sycl { +buffer buf{...}; + +void const_int_read_write(accessor acc) {...} +void int_read(accessor acc) {...} +void const_int_read(accessor acc) {...} + +q.submit([&](handler& cgh){ + auto accInt = + buf.get_access(cgh); // accessor + auto accIntConst = + accessor{buf, cgh}; // accessor + + // 1 + // accInt requested `read_write` mode, + // the scheduler may or may not be able to optimize this + const_int_read_write(accInt); + + // 2 + // Scheduler treats the combination of `const int` and `read_write` mode + // the same as if `read` mode was used, + // so this code is optimal + int_read(accIntConst); + + // 3 + // Similar to rule 2 + const_int_read(accIntConst); + + // 4 + // Allow all combinations + const_int_read_write(accIntConst); + const_int_read_write(accessor{}); + const_int_read_write(accessor{}); + int_read(accInt); + int_read(accessor{}); + int_read(accessor{}); + const_int_read(accInt); + const_int_read(accessor{}); + const_int_read(accessor{}); + + ... +}); +``` -template < - typename dataT, - int dims = 1, - access::mode accMode = access::mode::read_write, - access::target accTarget = access::target::global_buffer, - access::placeholder isPlaceholder = access::placeholder::false_t> -class accessor; +We also propose the following rules for simplifying the access mode: -} // namespace sycl -} // namespace cl +5. Convert an accessor of certain access modes + to an accessor of mode `access::mode::read_write`, + all other template parameters being equal + * Allowed modes are `access::mode::write`, `access::mode::discard_write`, `access::mode::discard_read_write` + +```cpp +buffer buf{...}; +q.submit([&](handler& cgh){ + // 5 + // The scheduler has been instructed to ignore previous data, + // even though the resulting accessor has a `read_write` mode, + // this is optimal code + accessor acc1 = buf.get_access(cgh); + + ... +}); ``` ## Aliases -We propose the following aliases to accessors: +With all of the above changes and simplifications applied, +we propose adding the following alias templates: ```cpp namespace cl { namespace sycl { -template -using buffer_accessor = - accessor::value, - access::target::global_buffer>; - template using constant_buffer_accessor = accessor; -template +template using host_accessor = accessor::value, + accMode, access::target::host_buffer>; } // namespace sycl } // namespace cl ``` -## Explicit conversions to aliases - -We propose allowing an accessor -that isn't quite the same type as one of the aliases - -for example, by having a `write` instead of a `read_write` access mode - -to be explicitly converted to an alias type. - -```cpp -namespace cl { -namespace sycl { - -template < - typename dataT, - int dims = 1, - access::mode accMode = access::mode::read_write, - access::target accTarget = access::target::global_buffer, - access::placeholder isPlaceholder = access::placeholder::false_t> -class accessor { - public: - /// All existing members here - - ... - - // Explicit conversion to `buffer_accessor` - // Only allowed when `accTarget == access::target::global_buffer` - explicit operator - buffer_accessor< - type_from_access_mode_t, dims>(); - - // Explicit conversion to `constant_buffer_accessor` - // Only allowed when `accTarget == access::target::constant_buffer` - explicit operator - constant_buffer_accessor< - type_from_access_mode_t, dims>(); - - // Explicit conversion to `host_accessor` - // Only allowed when `accTarget == access::target::host_buffer` - explicit operator - host_accessor< - type_from_access_mode_t, dims>(); -}; - -} // namespace sycl -} // namespace cl -``` - -| Member function | Description | -|-----------------|-------------| -| *`explicit operator buffer_accessor, dims>()`* | Performs a cast to `buffer_accessor`. Only allowed when `accTarget == access::target::global_buffer` and if `accMode` doesn't discard the constness of `dataT`. | -| *`explicit operator constant_buffer_accessor()`* | Performs a cast to `constant_buffer_accessor`. Only allowed when `accTarget == access::target::constant_buffer`. | -| *`explicit operator host_accessor, dims>()`* | Performs a cast to `host_accessor`. Only allowed when `accTarget == access::target::host_buffer` and if `accMode` doesn't discard the constness of `dataT`. | - -## Extending buffer class to return aliases - -In order to maintain backwards compatibility with SYCL 1.2.1, -the `get_access` member function should not change, -apart from defaulting the template parameters. -Instead, we propose adding new member functions to the `buffer` class -that request data access and return one of the new aliases. - -```cpp -namespace cl { -namespace sycl { - -template -class buffer { - public: - /// All existing members here - - ... - - - /// Existing `get_access` gains defaulted parameters - - template - accessor - get_access(handler& cgh); - - template - accessor - get_access(); - - - /// New functions - - template - buffer_accessor< - type_from_access_mode_t, dims> - get_device_access(handler& cgh); - - constant_buffer_accessor - get_device_constant_access(handler& cgh); +`constant_buffer_accessor` is very similar to `local_accessor` +in that it can only have one access mode. +It is allowed to use both `const dataT` and `dataT` as the data type, +the read-only access mode ensures data cannot be written to. - template - host_accessor< - type_from_access_mode_t, dims> - get_host_access(); -}; - -} // namespace sycl -} // namespace cl -``` - -| Member function | Description | -|-----------------|-------------| -| *`template buffer_accessor, dims> get_device_access(handler& cgh)`* | Calls `get_access`, but returns a `buffer_accessor`. | -| *`constant_buffer_accessor get_device_constant_access(handler& cgh)`* | Calls `get_access`, but returns a `constant_buffer_accessor`. | -| *`template host_accessor, dims> get_host_access()`* | Calls `get_access`, but returns a `host_accessor`. | +The `host_accessor` alias is similar to the regular global buffer accessor. ## Extending the handler @@ -329,13 +364,14 @@ which could now have less information to guide the scheduling process. This would be especially problematic for placeholder accessors, where the type of the accessor is determined before they are registered with a command group. -To resolve this, we propose extending the `handler` class -by allowing `require` to be called on any accessor -and to add an overload of `require` that also takes an access mode. -An example would be to call `require` with `global_accessor` -and `access::mode::discard_read_write`, -which would inform the scheduler to not copy over any old data. +To resolve this, we propose the following extensions to `handler::require`: + +* Allow it to be called on any non-host-mode accessor +* Add an overload that takes an access mode as a template parameter +* Return the accessor instance that was passed in + +See [Calling require](#calling-require) for an example. ```cpp namespace cl { @@ -349,22 +385,28 @@ class handler { /// New functions // Registers an accessor with a command group submission - template - void require(accessor acc); + accessor + require(accessor acc); // Registers an accessor with a command group submission - // `requestedMode` is a hint to the scheduler + // Already existed in 1.2.1, now it can take any accessor + // and return the same accessor back template - void require(accessor acc, - access::mode requestedMode); + accessor + require(accessor acc) { + return this->require(acc); + } }; } // namespace sycl @@ -373,8 +415,8 @@ class handler { | Member function | Description | |-----------------|-------------| -| *`template void require(accessor acc)`* | Registers the accessor for command group submission. | -| *`template void require(accessor acc, access::mode requestedMode)`* | Registers the accessor for command group submission. `requestedMode` is used as a hint to the scheduler. `requestedMode` cannot be a write mode if the accessor mode is `access::mode::read`. | +| *`template accessor require(accessor acc)`* | Registers the accessor for command group submission. Host accessors are not allowed. Returns `acc`. | +| *`template accessor require(accessor acc)`* | Registers the accessor for command group submission. `requestedMode` can be used to weaken the access mode. `requestedMode` cannot be a write mode if the accessor mode is `access::mode::read`. Host accessors are not allowed. Returns `acc`. | ## Considerations and alternatives @@ -400,8 +442,10 @@ in reducing verbosity there. ### Placeholder accessors -This proposal builds on top of allowing any accessor to be a placeholder. -The proposal still works without that assumption, +This proposal builds on top of accessors to be a placeholders without a template parameter: +[CP024 Default placeholders](https://github.com/codeplaysoftware/standards-proposals/pull/122). + +The current proposal still works without that assumption, but its usefulness is significantly reduced. If any accessor can be a placeholder, then the `isPlaceholder` template parameter becomes obsolete, @@ -412,6 +456,7 @@ placeholder accessors would not benefit from this proposal, unless some changes are made. Here are some options: + 1. Add an `isPlaceholder` parameter to the aliases, default it to `false_t`. This wouldn't reduce accessor verbosity as much as originally planned, by still requiring three template parameters for placeholder accessor aliases. @@ -419,58 +464,137 @@ Here are some options: We don't consider this option very desirable since it just replaces one kind of verbosity for another. -### Discouraging access mode combinations - -It is not clear to what extent should the specification discourage -the use of a non-const accessor `dataT` with `access::mode::read` -and a const accessor `dataT` with a write mode. -We considered deprecating those use cases, -or maybe just adding a warning, -or even `static_assert` to prevent their usage. -However, at least the example of `access::mode::read` with a non-const `dataT` -is very common in existing SYCL code -because using `const` is not considered essential in SYCL 1.2.1. +### Class template argument deduction + +It was suggested CTAD would help with some of the accessor verbosity +when compiling in C++17 mode. +We agree, but there are limitations: +C++17 doesn't allow CTAD on alias templates. +That would mean that while CTAD might work well with `access::target::global_buffer`, +since that's the default target and one can just use `accessor`, +it wouldn't work with `access::target::constant_buffer` +or `access::target::host_buffer` +since those rely on alias templates +`constant_buffer_accessor` and `host_accessor`, respectively. +This also affects the `local_accessor` alias already in SYCL 1.2.1. + +An option for solving this pre-C++20 would be +to define `constant_buffer_accessor` and `host_accessor` as new types +instead of alias templates. +They would publicly inherit from the `accessor` class +using the appropriate access target. + +However, this would also require defining additional constructors, +implicit conversions, and deduction guides for the feature to work as desired. + +### Detecting writeable accessors + +We considered adding a type trait that would indicate +whether the accessor is writeable or not. +This relates to [Constness of data](#constness-of-data) +and [Subscript operator](#subscript-operator), +where data cannot be modified if the underlying data type is `const`-qualified +or if the access mode is read-only. + +This could either be a `constexpr static` member of the accessor class +or a class template. ## Examples -### Storing accessors in a kernel functor +* [Simple global and host accessor use](#simple-global-and-host-accessor-use) +* [Vector addition](#vector-addition) +* [Simpler accessor construction](#simpler-accessor-construction) +* See [Storing accessors in a kernel functor](#constness-of-data) +* See [Implicit conversions](#implicit-conversions) -There is an example in the -[ComputeCpp SDK](https://github.com/codeplaysoftware/computecpp-sdk/blob/master/samples/template-function-object.cpp) -that showcases a kernel functor with stored accessors. -It uses user-defined aliases `read_accessor` and `write_accessor` -to simplify the three private members - -two read accessors and one write accessor. -These kinds of user-defined aliases seem to be a regular occurrence -in SYCL ecosystem code. -As shown by the example below, -this proposal would essentially standardize the different aliases -from across different projects. +### Simple global and host accessor use ```cpp -using namespace cl::sycl; +buffer buf{...}; +queue q; +q.submit([&](handler& cgh) { + accessor a{buf, cgh}; + ... // Write to buffer +}); +{ + host_accessor a{B}; + ... // Read from buffer +} +``` -template -class vector_add_kernel { - public: - vector_add_kernel(buffer_accessor ptrA, - buffer_accessor ptrB, - buffer_accessor ptrC) - : m_ptrA(ptrA), m_ptrB(ptrB), m_ptrC(ptrC) {} +### Vector addition - void operator()(item<1> item) { - m_ptrC[item.get_id()] = m_ptrA[item.get_id()] + m_ptrB[item.get_id()]; - } +The following example uses C++17 and CTAD. - private: - buffer_accessor m_ptrA; - buffer_accessor m_ptrB; - buffer_accessor m_ptrC; -}; +```cpp +std::vector a{1, 2, 3, 4, 5}; +std::vector b{6, 7, 8, 9, 10}; + +const auto N = a.size(); +const auto bufRange = range<1>(N); + +queue myQueue; + +// Create a buffer and copy `a` into it +buffer bufA{bufRange}; +myQueue.submit([&](handler &cgh) { + accessor accA{bufA, cgh}; // accessor + assert(!accA.is_null()); + assert(accA.has_handler()); + cgh.require( + accA); // Doesn't change type, scheduler can ignore previous data + cgh.copy(a.data(), accA); +}); + +// Create a buffer and copy `b` into it +// Use placeholders +accessor accB; // accessor +assert(accB.is_null()); +assert(!accB.has_handler()); +buffer bufB{bufRange}; +accB = accessor{bufB}; // accessor +assert(!accB.is_null()); +assert(!accB.has_handler()); +myQueue.submit([&](handler &cgh) { + cgh.require( + accB); // Doesn't change type, scheduler can ignore previous data + cgh.copy(b.data(), accB); +}); + +// Submit kernel that writes to output buffer +// Use constant buffer accessors +buffer bufC{bufRange}; +myQueue.submit([&](handler &cgh) { + accessor A{ + bufA, cgh}; // accessor + constant_buffer_accessor B{ + bufB, cgh}; // accessor + auto C = cgh.require( + accessor{bufC}); // accessor + cgh.parallel_for(bufRange, + [=](id<1> i) { C[i] = A[i] + B[i]; }); +}); +{ + // Request host access + host_accessor accC{ + bufC}; // accessor + assert(!accC.is_null()); + assert(!accC.has_handler()); + for (int i = 0; i < N; ++i) { + std::cout << accC[i] << std::endl; + } +} ``` -### New `get_access` functions +### Simpler accessor construction + +The following code shows two ways of expressing the same thing - +first in SYCL 1.2.1 code, and the second way according to this proposal. +It also assumes that accessor template parameters are defaulted in SYCL 1.2.1 +as proposed in +[Defaulting template parameters](#defaulting-template-parameters), +even though that's not allowed in SYCL 1.2.1. ```cpp using namespace cl::sycl; @@ -482,24 +606,22 @@ buffer buf; // Global buffer access accessor bufAcc = buf.get_access(cgh); -buffer_accessor bufAccNew = - buf.get_device_access(cgh); +accessor bufAccNew{buf, cgh}; // Global buffer access, read-only accessor bufAccRead = buf.get_access(cgh); -buffer_accessor bufAccReadNew = - buf.get_device_access(cgh); +accessor bufAccReadNew{buf, cgh}; // Global buffer access, ignore previous data accessor bufAccDiscard = buf.get_access(cgh); -buffer_accessor bufAccNewDiscard = - buf.get_device_access(cgh); +accessor bufAccNewDiscard = + buf.get_access(cgh); // Constant buffer access accessor bufAccConst = buf.get_access(cgh); -constant_buffer_accessor bufAccConstNew = - buf.get_device_constant_access(cgh); +constant_buffer_accessor bufAccConstNew{buf, cgh}; // Host buffer access accessor bufAccHost = buf.get_access(); -host_accessor bufAccHostNew = - buf.get_host_access(); +host_accessor bufAccHostNew{buf}; // Host buffer access, read-only accessor bufAccHostRead = buf.get_access(); -host_accessor bufAccHostNewRead = - buf.get_host_access(); +host_accessor bufAccHostNewRead{buf}; // Host buffer access, ignore previous data accessor bufAccHostDiscard = buf.get_access(); host_accessor bufAccHostNewDiscard = - buf.get_host_access(); + buf.get_access(); ``` ### Calling `require` @@ -545,27 +664,30 @@ using namespace cl::sycl; // Assume these are available queue q; -buffer buf1; -buffer buf2; +buffer bufA; +buffer bufB; +buffer bufC; q.submit([](handler& cgh) { - // Request read-write access to both buffers - buffer_accessor bufAcc1 = - buf1.get_device_access(cgh); - buffer_accessor bufAcc2 = - buf2.get_device_access(cgh); - - // Register the accessor for command group submission + // Request read-only access to bufA + accessor accA{bufA, cgh}; + // Register accA for command group submission // Not so useful in this case since it's already been registered, - // but bufAcc1 could also be a placeholder - cgh.require(bufAcc1); + // but accA could also be a placeholder + cgh.require(accA); - // Register the accessor for command group submission - // bufAcc2 has also already been registered, - // but this instructs the scheduler to ignore previous data - cgh.require(bufAcc2, access::mode::discard_read_write); + // Create a placeholder with read-only access to bufB + // The accessor is immediately registered and returned + // accB is of type accessor + auto accB = cgh.require(accessor{bufB}); + + // Create a placeholder with read-write access to bufC + // The accessor is immediately registered and returned + // The provided access mode instructs the scheduler to ignore previous data + // and "weaken" the scheduling mode to write-only + // accC is of type accessor + auto accC = cgh.require(accessor{bufC}); ... }); - ``` From 7905e94487a6ecf711e16d907132de1307e78266 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Peter=20=C5=BDu=C5=BEek?= Date: Sun, 5 Apr 2020 21:19:02 +0100 Subject: [PATCH 3/3] CP023 Accessor aliases, revision 0.3 In order to reduce the verbosity of programming with SYCL accessors, this proposal aims to reduce the number of template parameters to just 2 from the current 5. It achieves this by slightly revising how to treat read-only data and adding alias templates based on the access target. 1. Main changes * Default accessor template parameters * Simplify access modes * Alias templates based on access target * Extend the handler class 2. Accessor alias templates * `constant_buffer_accessor * `host_accessor` as a new type 3. Treat `const T` the same as `access::mode::read` * Simplifies a lot of code 4. Define implicit conversions for equivalent types 5. Implicit conversions that add `const` 6. Added `property::discard` 7. Overload for `handler::require` that also takes `property::discard` * To ignore previous data 8. Default all accessor template parameters * Expect the data type * Default to an accessor to global buffer with read-write access * Assumes accessors can be a placeholder without template parameter 9. Good support for CTAD * Deduction tags as compile-time properties * Made all accessor constructors variadic templates 10. Discussed some considerations and alternatives 11. Examples of reduced verbosity --- README.md | 2 +- accessor-alias/index.md | 8 +- accessor-alias/sycl-2.2/index.md | 259 +++++++++++++++++++++++-------- 3 files changed, 198 insertions(+), 71 deletions(-) diff --git a/README.md b/README.md index aa3dfc0..61171c9 100644 --- a/README.md +++ b/README.md @@ -63,5 +63,5 @@ from this registry in the future. | CP020 | [Interop Task](interop_task/interop_task.md) | SYCL 1.2.1 | 16 January 2019 | 16 January 2019 | _Available since CE 1.0.5_ | | CP021 | [Default-Constructed Buffers](default-constructed-buffers/default-constructed-buffers.md) | SYCL 1.2.1 | 27 August 2019 | 5 September 2019 | _Draft_ | | CP022 | [Host Task with Interop capabilities](host_task/host_task.md) | SYCL 1.2.1 | 16 January 2019 | 20 January 2020 | _Final Draft_ | -| CP023 | [Accessor aliases](accessor-alias/index.md) | SYCL Next (after 1.2.1) | 22 September 2019 | 31 March 2020 | _Work in Progress_ | +| CP023 | [Accessor aliases](accessor-alias/index.md) | SYCL Next (after 1.2.1) | 22 September 2019 | 6 April 2020 | _Work in Progress_ | | CP026 | [Generalized Error Handling For SYCL](error-handling/sycl-error-handling.md) | SYCL Next | 10 March 2020 | 10 March 2020 | _Under Review_ | diff --git a/accessor-alias/index.md b/accessor-alias/index.md index 8811e5e..2ab4473 100644 --- a/accessor-alias/index.md +++ b/accessor-alias/index.md @@ -4,13 +4,13 @@ |-------------|--------| | Name | Accessor aliases | | Date of Creation | 22 September 2019 | -| Revision | 0.2 | -| Latest Update | 31 March 2020 | +| Revision | 0.3 | +| Latest Update | 6 April 2020 | | Target | SYCL Next (after 1.2.1) | | Current Status | _Work in Progress_ | | Reply-to | Peter Žužek | | Original author | Peter Žužek | -| Contributors | | +| Contributors | Gordon Brown gordon@codeplay.com | ## Overview @@ -19,4 +19,4 @@ in order to reduce accessor verbosity. ## Versions -[Revision 0.2](sycl-2.2/index.md) +[Revision 0.3](sycl-2.2/index.md) diff --git a/accessor-alias/sycl-2.2/index.md b/accessor-alias/sycl-2.2/index.md index fefbb08..4a44871 100644 --- a/accessor-alias/sycl-2.2/index.md +++ b/accessor-alias/sycl-2.2/index.md @@ -45,6 +45,20 @@ local_accessor ## Revisions +### 0.3 + +* Support CTAD a lot better, including in examples +* Consider compile-time and run-time properties +* Added deduction tags + * Considered compile-time properties + * `read_only_tag` deduces `const T` and `access::mode::read` + * `constant_access_tag` deduces accessor to constant memory +* Added `property::discard` +* `handler::require` doesn't take extra template parameter anymore, + instead it can accept `property::discard` as an argument +* Made `host_accessor` a separate type +* Use `empty()` instead of `is_null()` + ### 0.2 * Simplified handling of non-writeable accessors @@ -75,11 +89,15 @@ Main points of the proposal: * `read` and `read_write` most important * Allow `const T` to denote read-only data access * Deprecate `access::mode::atomic` +1. [Introduce deduction tags](#deduction-tags) +1. [Extend properties](#extending-properties) 1. [Introduce `access::target`-specific aliases for accessors](#aliases) 1. [Extending `handler::require`](#extending-the-handler) The proposal slightly changes the semantics of accessing data, but still aims to be completely backwards compatible with SYCL 1.2.1. +There is a lot of focus on CTAD which is a C++17 feature, +but the proposal can still work with C++11 if CTAD is ignored. Note that this proposal assumes that any accessor can be a placeholder, essentially deprecating the `access::placeholder` template parameter. @@ -231,7 +249,7 @@ In order to simplify user code, we propose allowing certain implicit conversions that don't modify scheduling information. -In standard C++ code, adding `const` qualifiers is almost always allowed. +In standard C++ code, adding the `const` qualifier is almost always allowed. In SYCL, going from a `read_write` accessor to a `read` accessor is analogous to adding the `const` qualifier. This allows us to have the following rules regarding `const`: @@ -320,15 +338,137 @@ q.submit([&](handler& cgh){ }); ``` +## Deduction tags + +In order to better enable CTAD, +we propose the following deduction tags: + +```cpp +namespace cl { +namespace sycl { + +struct read_only_tag {}; +struct constant_access_tag {}; + +} // namespace sycl +} // namespace cl +``` + +`read_only_tag` can be used to aid deduction guides +in deducing read-only access to an accessor: +it deduces the accessor to `const dataT` data type and `access::mode::read`. + +`constant_access_tag` can be used to aid deduction guides +in deducing access to constant buffer memory: +it deduces the accessor to `const dataT` data type, +`access::mode::read`, +and `access::target::constant_buffer`. +It can be used in conjunction with `read_only_tag`. + +When used in a non-CTAD context, +deduction tags have no effect, +but the implementation must reject code +that would lead to incompatible accessor types. + +## Extending properties + +### Compile-time properties + +SYCL 1.2.1 already supports passing run-time properties to accessors +(see https://github.com/KhronosGroup/SYCL-Docs/pull/73). + +Consider the following SYCL 1.2.1 constructor +for `accessor`: + +```cpp +accessor(buffer&, handler&, property_list = {}); +``` + +It is possible to construct an object of `property_list` +with accessor specific properties and pass it to the accessor constructor. + +We propose treating deduction tags as compile-time properties. +Compile-time properties are not passed to `property_list` - +that class still encapsulates run-time properties. +Instead, compile-time properties need to help deduction guides, +which requires changing all accessor constructors +to variadic templates. +The example constructor above would thus become: + +```cpp +template +accessor(buffer&, handler&, PropertyTs...); +``` + +Passing compile-time properties to `property_list` has no effect +and they can be ignored, unless specified otherwise. +But this variadic constructor allows us to pass in deduction tags to guide CTAD. +The following would be valid: + +```cpp +// Assume buffer `buf` +// Assume command group handler `cgh` +// Assume property_list `propList` +accessor acc{buf, cgh, read_only_tag{}, constant_access_tag{}, propList}; +static_assert(std::is_same_v< + decltype(acc), + accessor>); +``` + +This example also demonstrates the only allowed order of all arguments: +deduction tags come after the main arguments, +and before the run-time `property_list`. +However, any property can be omitted. + +> Note that `read_only_tag` is not required +> when `constant_access_tag` is specified. + +### Discard property + +We propose a new run-time property that can be passed to accessor constructors +inside `property_list`: + +```cpp +namespace cl { +namespace sycl { +namespace property { + +struct discard {}; +__unspecified__ constexpr auto discard_v = discard{}; + +} // namespace property +} // namespace sycl +} // namespace cl +``` + +This property is a signal to the scheduler +that the previous data should be ignored. + ## Aliases With all of the above changes and simplifications applied, -we propose adding the following alias templates: +we propose adding the following new classes and alias templates: ```cpp namespace cl { namespace sycl { +template +class host_accessor : + public accessor { + public: + /// Constructors + ... + + /// Implicit conversions + ... +}; + template using constant_buffer_accessor = accessor; -template -using host_accessor = - accessor; - } // namespace sycl } // namespace cl ``` +`host_accessor` is a new type because it allows the use of CTAD. +Its base type is `accessor`. +`host_accessor` is implicitly convertible to and from the base type. + `constant_buffer_accessor` is very similar to `local_accessor` in that it can only have one access mode. It is allowed to use both `const dataT` and `dataT` as the data type, the read-only access mode ensures data cannot be written to. -The `host_accessor` alias is similar to the regular global buffer accessor. - ## Extending the handler Reducing effective access modes from 6 to 2 @@ -368,7 +501,7 @@ before they are registered with a command group. To resolve this, we propose the following extensions to `handler::require`: * Allow it to be called on any non-host-mode accessor -* Add an overload that takes an access mode as a template parameter +* Add an overload that takes `property::discard` as an argument * Return the accessor instance that was passed in See [Calling require](#calling-require) for an example. @@ -385,9 +518,9 @@ class handler { /// New functions // Registers an accessor with a command group submission - // `requestedMode` can be used to weaken the access mode - template acc); // Registers an accessor with a command group submission - // Already existed in 1.2.1, now it can take any accessor - // and return the same accessor back + // Overload that instructs the scheduler to discard previous data template accessor - require(accessor acc) { - return this->require(acc); - } + require(accessor acc, + property::discard); }; } // namespace sycl @@ -416,7 +547,7 @@ class handler { | Member function | Description | |-----------------|-------------| | *`template accessor require(accessor acc)`* | Registers the accessor for command group submission. Host accessors are not allowed. Returns `acc`. | -| *`template accessor require(accessor acc)`* | Registers the accessor for command group submission. `requestedMode` can be used to weaken the access mode. `requestedMode` cannot be a write mode if the accessor mode is `access::mode::read`. Host accessors are not allowed. Returns `acc`. | +| *`template accessor require(accessor acc, property::discard)`* | Registers the accessor for command group submission and instructs the scheduler to ignore previous data. Host accessors are not allowed. Returns `acc`. | ## Considerations and alternatives @@ -464,29 +595,6 @@ Here are some options: We don't consider this option very desirable since it just replaces one kind of verbosity for another. -### Class template argument deduction - -It was suggested CTAD would help with some of the accessor verbosity -when compiling in C++17 mode. -We agree, but there are limitations: -C++17 doesn't allow CTAD on alias templates. -That would mean that while CTAD might work well with `access::target::global_buffer`, -since that's the default target and one can just use `accessor`, -it wouldn't work with `access::target::constant_buffer` -or `access::target::host_buffer` -since those rely on alias templates -`constant_buffer_accessor` and `host_accessor`, respectively. -This also affects the `local_accessor` alias already in SYCL 1.2.1. - -An option for solving this pre-C++20 would be -to define `constant_buffer_accessor` and `host_accessor` as new types -instead of alias templates. -They would publicly inherit from the `accessor` class -using the appropriate access target. - -However, this would also require defining additional constructors, -implicit conversions, and deduction guides for the feature to work as desired. - ### Detecting writeable accessors We considered adding a type trait that would indicate @@ -513,11 +621,11 @@ or a class template. buffer buf{...}; queue q; q.submit([&](handler& cgh) { - accessor a{buf, cgh}; + accessor a{buf, cgh}; ... // Write to buffer }); { - host_accessor a{B}; + host_accessor a{B, read_only_tag{}}; ... // Read from buffer } ``` @@ -539,25 +647,27 @@ queue myQueue; buffer bufA{bufRange}; myQueue.submit([&](handler &cgh) { accessor accA{bufA, cgh}; // accessor - assert(!accA.is_null()); + assert(!accA.empty()()); assert(accA.has_handler()); - cgh.require( - accA); // Doesn't change type, scheduler can ignore previous data + cgh.require( + accA, property::discard_v); // Doesn't change type, + // scheduler can ignore previous data cgh.copy(a.data(), accA); }); // Create a buffer and copy `b` into it // Use placeholders accessor accB; // accessor -assert(accB.is_null()); +assert(accB.empty()()); assert(!accB.has_handler()); buffer bufB{bufRange}; accB = accessor{bufB}; // accessor -assert(!accB.is_null()); +assert(!accB.empty()()); assert(!accB.has_handler()); myQueue.submit([&](handler &cgh) { - cgh.require( - accB); // Doesn't change type, scheduler can ignore previous data + cgh.require( + accB, property::discard_v); // Doesn't change type, + // scheduler can ignore previous data cgh.copy(b.data(), accB); }); @@ -565,12 +675,20 @@ myQueue.submit([&](handler &cgh) { // Use constant buffer accessors buffer bufC{bufRange}; myQueue.submit([&](handler &cgh) { - accessor A{ - bufA, cgh}; // accessor - constant_buffer_accessor B{ - bufB, cgh}; // accessor - auto C = cgh.require( - accessor{bufC}); // accessor + accessor A{bufA, cgh, read_only_tag{}}; + static_assert(std::is_same_v< + decltype(A), + accessor>); + accessor B{bufA, cgh, constant_access_tag{}}; + static_assert(std::is_same_v< + decltype(B), + accessor>); + auto C = cgh.require(accessor{bufC}, property::discard_v); + static_assert(std::is_same_v< + decltype(C), + accessor>); cgh.parallel_for(bufRange, [=](id<1> i) { C[i] = A[i] + B[i]; }); }); @@ -579,7 +697,7 @@ myQueue.submit([&](handler &cgh) { // Request host access host_accessor accC{ bufC}; // accessor - assert(!accC.is_null()); + assert(!accC.empty()()); assert(!accC.has_handler()); for (int i = 0; i < N; ++i) { std::cout << accC[i] << std::endl; @@ -607,6 +725,7 @@ buffer buf; accessor bufAcc = buf.get_access(cgh); accessor bufAccNew{buf, cgh}; +accessor bufAccCTAD{buf, cgh}; // Global buffer access, read-only accessor bufAccRead = buf.get_access(cgh); accessor bufAccReadNew{buf, cgh}; +accessor bufAccReadCTAD{buf, cgh, read_only_tag{}}; // Global buffer access, ignore previous data accessor(cgh); accessor bufAccNewDiscard = buf.get_access(cgh); +auto bufAccNewDiscard2 = accessor{buf, cgh, {property::discard_v}}; +accessor bufAccCTADDiscard{buf, cgh, {property::discard_v}}; // Constant buffer access accessor bufAccConst = buf.get_access(cgh); constant_buffer_accessor bufAccConstNew{buf, cgh}; +accessor bufAccConstCTAD{buf, cgh, constant_access_tag{}}; // Host buffer access accessor bufAccHost = buf.get_access(); host_accessor bufAccHostNew{buf}; +host_accessor bufAccHostCTAD{buf}; // Host buffer access, read-only accessor bufAccHostRead = buf.get_access(); host_accessor bufAccHostNewRead{buf}; +host_accessor bufAccHostCTADRead{buf, read_only_tag{}}; // Host buffer access, ignore previous data accessor(); host_accessor bufAccHostNewDiscard = buf.get_access(); +auto bufAccHostNewDiscard2 = host_accessor{buf, {property::discard_v}}; +host_accessor bufAccHostCTADDiscard{buf, {property::discard_v}}; ``` ### Calling `require` @@ -670,7 +797,7 @@ buffer bufC; q.submit([](handler& cgh) { // Request read-only access to bufA - accessor accA{bufA, cgh}; + accessor accA{bufA, cgh, read_only_tag{}}; // Register accA for command group submission // Not so useful in this case since it's already been registered, // but accA could also be a placeholder @@ -683,10 +810,10 @@ q.submit([](handler& cgh) { // Create a placeholder with read-write access to bufC // The accessor is immediately registered and returned - // The provided access mode instructs the scheduler to ignore previous data - // and "weaken" the scheduling mode to write-only + // The provided discard property instructs the scheduler + // to ignore previous data // accC is of type accessor - auto accC = cgh.require(accessor{bufC}); + auto accC = cgh.require(accessor{bufC}, property::discard_v); ... });