diff --git a/README.md b/README.md index 25db876..20c8cf6 100644 --- a/README.md +++ b/README.md @@ -63,4 +63,4 @@ 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 | 9 March 2020 | _Work in Progress_ | diff --git a/accessor-alias/index.md b/accessor-alias/index.md index 8a6615f..a69753e 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 | 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 | @@ -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..fd02d71 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,180 @@ 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) + int_read(accInt) + const_int_read(accInt) + + ... +}); +``` -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 +`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. -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`. | +The `host_accessor` alias is similar to the regular global buffer accessor. ## Extending the handler @@ -329,13 +358,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 +379,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 +409,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 +436,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 +450,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 +458,136 @@ 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 +### 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. -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. +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 +599,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 +657,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}); ... }); - ```