From 120b04786c1c8a08e7120be4c0cbe7e65b741a20 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Peter=20=C5=BDu=C5=BEek?= Date: Mon, 9 Mar 2020 08:13:14 +0000 Subject: [PATCH 1/4] CP024 Default placeholders 1. Deprecate `access::placeholder` 2. Extend placeholder functionality to ignore template parameter * Still applies only to global and constant buffers 3. Updated `handler::require` 4. New accessor constructors * Default constructor * From a buffer, 0-dim * From a buffer, ranged 5. New access member functions * `is_null` * `has_handler` 6. Vector addition example --- README.md | 3 +- default-placeholders/index.md | 269 ++++++++++++++++++++++++++++++++++ 2 files changed, 271 insertions(+), 1 deletion(-) create mode 100644 default-placeholders/index.md diff --git a/README.md b/README.md index 5099833..ff658f6 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 +| CP024 | [Default placeholder](default-placeholders/index.md) | SYCL Next (after 1.2.1) | 9 March 2020 | 9 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/default-placeholders/index.md b/default-placeholders/index.md new file mode 100644 index 0000000..b4ebbb8 --- /dev/null +++ b/default-placeholders/index.md @@ -0,0 +1,269 @@ +# Default placeholders + +| Proposal ID | CP024 | +|-------------|--------| +| Name | Default placeholders | +| Date of Creation | 9 March 2019 | +| Revision | 0.1 | +| Latest Update | 9 March 2020 | +| Target | SYCL Next (after 1.2.1) | +| Current Status | _Work in Progress_ | +| Reply-to | Peter Žužek | +| Original author | Peter Žužek | +| Contributors | Gordon Brown | + +## Overview + +This proposal aims to deprecate `access::placeholder` +and instead allow all accessors to global and constant memory to be placeholders. + +## Revisions + +### 0.1 + +* Initial proposal + +## Motivation + +SYCL 1.2.1 introduced the `access::placeholder` enumeration +which is used as the 5th accessor template parameter +to indicate whether the accessor can be used as a placeholder. +Only the `global_buffer` and `constant_buffer` access targets +support placeholder accessors. + +The main reason for having placeholders is to store accessors into objects +without having a queue at the point of object creation, +and registering the access with a command group submission at a later time. +One of Codeplay's proposals that didn't make it into SYCL 1.2.1 +was to allow placeholders to be default constructible as well, +i.e. not having to even know the buffer at the point of construction. +This extension is used in some SYCL ecosystem projects. + +Accessors are pointer-like objects, +and placeholders try to fill that gap +that prevents accessors from being even more pointer-like. +A default constructed placeholder accessor is analogous to a null pointer. +A placeholder that's bound to a buffer +but hasn't been registered with a command group +is more like a fancy pointer: +the user doesn't own the data +until the accessor is registered and used in a kernel, +where is becomes more similar to a regular pointer. + +Having this type separation between full accessors and placeholders +might be useful from a type safety perspective, +but we believe it makes development more difficult. +For example, [another one of our proposals](https://github.com/codeplaysoftware/standards-proposals/pull/100/files) +introduces alias templates for different access targets +and revises rules on how read-only accessors are handled, +all in the name of reducing accessor verbosity. +The placeholder template parameter makes that much more difficult, +meaning we either need to introduce another parameter to the alias template, +making it a lot less useful, +or simply ignore the reduction in verbosity for placeholder accessors. + +## Changes + +### Deprecate `access::placeholder` + +Mark the `access::placeholder` enum class as deprecated, +but keep it for backwards compatibility +until it's eventually removed from a subsequent standard. + +```cpp +namespace access { +... + +enum class placeholder // deprecated +{...}; +} // namespace access + +template +class accessor; +``` + +### All accessors with a global or constant target can be placeholders + +SYCL 1.2.1 allows `access::placeholder::true_t` +only when the access target is `global_buffer` or `constant_buffer`. +We propose that the same placeholder semantics +still apply to only these two targets, +just that the semantics and API is available +regardless of the `isPlaceholder` template parameter. + +### Accept all accessors in `handler::require` + +At the moment the member function `handler::require` only accepts +placeholder accessors. +This should stay the same, +but since the enum class is deprecated, +the function needs to be extended: + +```cpp +class handler { + public: + ... + + // Adds isPlaceholder to existing SYCL 1.2.1 function + template + void require(accessor + acc); +}; +``` + +### Deprecate `is_placeholder` + +The function `accessor::is_placeholder` doesn't make sense anymore, +we propose deprecating it. + +### New constructors + +We propose adding new constructors to the accessors class +to allow placeholder construction. + +1. Default constructor - not part of SYCL 1.2.1, + but there is a [Codeplay proposal](https://github.com/codeplaysoftware/standards-proposals/pull/89) + for making that happen. +1. 0-dim accessor constructor from a buffer - + normally an accessor constructor requires a buffer and a handler, + we propose making the handler optional. + This is the same constructor currently allowed for host buffers. +1. Constructor from a buffer - + normally an accessor constructor requires a buffer and a handler, + we propose making the handler optional. + This is the same constructor currently allowed for host buffers. + +```cpp +template +class accessor { + public: + ... + + // 1 + // Only available when ((accessTarget == access::target::global_buffer) || + // (accessTarget == access::target::constant_buffer)) + accessor() noexcept; + + // 2 + // Only available when: ((accessTarget == access::target::global_buffer) || + // (accessTarget == access::target::constant_buffer) || + // (accessTarget == access::target::host_buffer)) && + // (dimensions == 0) + accessor(buffer &bufferRef); + + // 3 + // Only available when: ((accessTarget == access::target::global_buffer) || + // (accessTarget == access::target::constant_buffer) || + // (accessTarget == access::target::host_buffer)) && + // (dimensions > 0) + accessor(buffer &bufferRef, + range accessRange, + id accessOffset = {}); +}; +``` + +### New `accessor` member functions + +In order to query the accessor for its status, +we propose new member functions to the `accessor class`: + +1. `is_null` - returns `true` if the accessor has been default constructed, + which is only possible with placeholders. + Not having an associated buffer is analogous to a null pointer. +1. `has_handler` - returns `true` if the accessor is associated + with a command group `handler`. + Will only be `false` with host accessors and placeholder accessors. + This replaces the `is_placeholder` member function. + +```cpp +template +class accessor { + public: + ... + + // 1 + bool is_null() const noexcept; + + // 2 + bool has_handler() const noexcept; +}; +``` + +## Examples + +Simple vector addition: + +```cpp +std::vector a{1, 2, 3, 4, 5}; +std::vector b{6, 7, 8, 9, 10}; + +using read_acc = + accessor; +using write_acc = + accessor; + +read_acc accA; +read_acc accB; +write_acc accC; + +// Sanity checks +assert(accA.is_null()); +assert(!accA.has_handler()); + +const auto N = a.size(); +const auto bufRange = range<1>(N); + +queue myQueue; + +// Create a buffer and copy `a` into it +buffer bufA{bufRange}; + +accA = read_acc{bufA}; +assert(!accA.is_null()); +assert(!accA.has_handler()); + +myQueue.submit([&](handler &cgh) { + cgh.require(accA); + cgh.copy(a.data(), accA); +}); + +// Create a buffer and copy `b` into it +buffer bufB{bufRange}; +accB = read_acc{bufB}; +myQueue.submit([&](handler &cgh) { + cgh.require(accB); + cgh.copy(b.data(), accB); +}); + +// Submit kernel that writes to output buffer +// Use constant buffer accessors +buffer bufC{bufRange}; +accC = read_acc{bufC}; +myQueue.submit([&](handler &cgh) { + cgh.require(accA); + cgh.require(accB); + cgh.require(accC); + cgh.parallel_for(bufRange, + [=](id<1> i) { accC[i] = accA[i] + accB[i]; }); +}); +``` From 191abb8b0f781eff8853bf2b36872260666eda60 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Peter=20=C5=BDu=C5=BEek?= Date: Fri, 13 Mar 2020 09:41:35 +0000 Subject: [PATCH 2/4] MR Feedback * Null accessors are allowed in kernels * But cannot be registered with a handler * Allow constructing host accessors from placeholders * Accessor constructor * `accessor::get_host_access` --- README.md | 2 +- default-placeholders/index.md | 84 +++++++++++++++++++++++++++++++---- 2 files changed, 77 insertions(+), 9 deletions(-) diff --git a/README.md b/README.md index ff658f6..59be65e 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_ | -| CP024 | [Default placeholder](default-placeholders/index.md) | SYCL Next (after 1.2.1) | 9 March 2020 | 9 March 2020 | _Work in Progress_ | +| CP024 | [Default placeholder](default-placeholders/index.md) | SYCL Next (after 1.2.1) | 9 March 2020 | 17 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/default-placeholders/index.md b/default-placeholders/index.md index b4ebbb8..c6529fc 100644 --- a/default-placeholders/index.md +++ b/default-placeholders/index.md @@ -4,8 +4,8 @@ |-------------|--------| | Name | Default placeholders | | Date of Creation | 9 March 2019 | -| Revision | 0.1 | -| Latest Update | 9 March 2020 | +| Revision | 0.2 | +| Latest Update | 17 March 2020 | | Target | SYCL Next (after 1.2.1) | | Current Status | _Work in Progress_ | | Reply-to | Peter Žužek | @@ -48,7 +48,7 @@ but hasn't been registered with a command group is more like a fancy pointer: the user doesn't own the data until the accessor is registered and used in a kernel, -where is becomes more similar to a regular pointer. +where it becomes more similar to a regular pointer. Having this type separation between full accessors and placeholders might be useful from a type safety perspective, @@ -123,11 +123,43 @@ class handler { }; ``` +`handler::require` has to be called on a placeholder accessor +in order to register it with the command group submission. +It is valid to call the function more than once, +even on non-placeholder accessors. +Calling the function on a null accessor throws `cl::sycl::invalid_object_error`. + ### Deprecate `is_placeholder` The function `accessor::is_placeholder` doesn't make sense anymore, we propose deprecating it. +### Allow constructing host accessors from placeholders + +Consider the following example: + +```cpp +template +void some_library_function(AccTypeA accA, AccTypeB accB) { + ... + myQueue.submit([&](handler &cgh) { + cgh.require(accA); + cgh.require(accB); + cgh.copy(accA, accB); + }); + ... + // We want to be able to access host data now +} +``` + +`some_library_function` in the example takes in two placeholder accessors +and performs a copy from one to another. +However, there is no way any of the data associated with the accessors +can be accessed on the host. +The placeholders are not bound to a command group anyway, +so we believe it should be possible to explicitly construct a host accessor +from a placeholder accessor. + ### New constructors We propose adding new constructors to the accessors class @@ -144,6 +176,17 @@ to allow placeholder construction. normally an accessor constructor requires a buffer and a handler, we propose making the handler optional. This is the same constructor currently allowed for host buffers. +1. Construct a host accessor from a placeholder one (`placeholderAcc`). + Not valid to call in kernel code. + Throws `cl::sycl::runtime_error` when called + if `placeholderAcc.has_handler() == true`. + Requesting host access is a synchronization point, + and host accessors act as locks, + meaning that the placeholder cannot be used + while the host accessor is in scope. + Even after host access is released, + the programmer is required to call `require` again on the placeholder + before it can be used in a kernel. ```cpp template &bufferRef); // 3 // Only available when: ((accessTarget == access::target::global_buffer) || - // (accessTarget == access::target::constant_buffer) || - // (accessTarget == access::target::host_buffer)) && + // (accessTarget == access::target::constant_buffer)) && // (dimensions > 0) accessor(buffer &bufferRef, range accessRange, id accessOffset = {}); + + // 4 + // Only available when (accessTarget == access::target::host_buffer) && + // ((otherTarget == access::target::global_buffer) || + // (otherTarget == access::target::constant_buffer)) + template + accessor(accessor& + placeholderAcc); }; ``` @@ -186,10 +239,16 @@ we propose new member functions to the `accessor class`: 1. `is_null` - returns `true` if the accessor has been default constructed, which is only possible with placeholders. Not having an associated buffer is analogous to a null pointer. + Available in both application code and kernel code, + it is valid to pass a null accessor to a kernel. 1. `has_handler` - returns `true` if the accessor is associated with a command group `handler`. Will only be `false` with host accessors and placeholder accessors. This replaces the `is_placeholder` member function. + Mainly meant as a way to enquire about whether this is a placeholder or not, + this doesn't have to be checked before `require` is called. +1. `get_host_access` - constructs a host accessor from a placeholder accessor. + Not valid to call in kernel code. ```cpp template + get_host_access() const; }; ``` @@ -256,7 +325,6 @@ myQueue.submit([&](handler &cgh) { }); // Submit kernel that writes to output buffer -// Use constant buffer accessors buffer bufC{bufRange}; accC = read_acc{bufC}; myQueue.submit([&](handler &cgh) { From fe8488ce3aac84ce09c1d96a0a8cb11bcfb0272c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Peter=20=C5=BDu=C5=BEek?= Date: Tue, 31 Mar 2020 10:21:10 +0100 Subject: [PATCH 3/4] Made constructors explicit --- README.md | 2 +- default-placeholders/index.md | 4 +++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/README.md b/README.md index 59be65e..399bcc5 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_ | -| CP024 | [Default placeholder](default-placeholders/index.md) | SYCL Next (after 1.2.1) | 9 March 2020 | 17 March 2020 | _Work in Progress_ | +| CP024 | [Default placeholder](default-placeholders/index.md) | SYCL Next (after 1.2.1) | 9 March 2020 | 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/default-placeholders/index.md b/default-placeholders/index.md index c6529fc..dabc29e 100644 --- a/default-placeholders/index.md +++ b/default-placeholders/index.md @@ -5,7 +5,7 @@ | Name | Default placeholders | | Date of Creation | 9 March 2019 | | Revision | 0.2 | -| Latest Update | 17 March 2020 | +| Latest Update | 31 March 2020 | | Target | SYCL Next (after 1.2.1) | | Current Status | _Work in Progress_ | | Reply-to | Peter Žužek | @@ -207,6 +207,7 @@ class accessor { // Only available when: ((accessTarget == access::target::global_buffer) || // (accessTarget == access::target::constant_buffer)) && // (dimensions == 0) + explicit accessor(buffer &bufferRef); // 3 @@ -222,6 +223,7 @@ class accessor { // ((otherTarget == access::target::global_buffer) || // (otherTarget == access::target::constant_buffer)) template + explicit accessor(accessor Date: Tue, 31 Mar 2020 10:25:43 +0100 Subject: [PATCH 4/4] Clarified `has_handler` --- default-placeholders/index.md | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/default-placeholders/index.md b/default-placeholders/index.md index dabc29e..319e637 100644 --- a/default-placeholders/index.md +++ b/default-placeholders/index.md @@ -129,6 +129,9 @@ It is valid to call the function more than once, even on non-placeholder accessors. Calling the function on a null accessor throws `cl::sycl::invalid_object_error`. +Note that `handler::require` does not attach a handler to an accessor, +which means it doesn't change the outcome of `accessor::has_handler`. + ### Deprecate `is_placeholder` The function `accessor::is_placeholder` doesn't make sense anymore, @@ -243,12 +246,13 @@ we propose new member functions to the `accessor class`: Not having an associated buffer is analogous to a null pointer. Available in both application code and kernel code, it is valid to pass a null accessor to a kernel. -1. `has_handler` - returns `true` if the accessor is associated +1. `has_handler` - returns `true` if the accessor was constructed with a command group `handler`. Will only be `false` with host accessors and placeholder accessors. This replaces the `is_placeholder` member function. Mainly meant as a way to enquire about whether this is a placeholder or not, this doesn't have to be checked before `require` is called. + Calling `require` does not change the result of `has_handler`. 1. `get_host_access` - constructs a host accessor from a placeholder accessor. Not valid to call in kernel code.