diff --git a/README.md b/README.md index 5099833..d92ff3d 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 +| CP026 | [Generalized Error Handling For SYCL](error-handling/sycl-error-handling.md) | SYCL Next | 10 March 2020 | 10 March 2020 | _Under Review_ | +| CP027 | [Accessors as Containers](accessor-container/index.md) | SYCL Next | 22 March 2020 | 22 March 2020 | _Work in Progress_ | diff --git a/accessor-container/index.md b/accessor-container/index.md new file mode 100644 index 0000000..699777d --- /dev/null +++ b/accessor-container/index.md @@ -0,0 +1,109 @@ +# Accessors as Containers + +| Proposal ID | CP027 | +|-------------|--------| +| Name | Accessors as Containers | +| Date of Creation | 22 March 2019 | +| Revision | 0.1 | +| Latest Update | 22 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 introduce a `Container` interface for accessors. + +## Revisions + +### 0.1 + +* Initial proposal + +## Motivation + +Users want to be able to iterate over accessors +and pass them to algorithms that rely on iterators. + +## Changes + +Apply the changes to accessors of the following access targets: + +* `global_buffer` +* `constant_buffer` +* `local` +* `host_buffer` + +Add a section to the spec that says roughly the following: +> Accessors of certain access targets +meet the C++ ContiguousContainer named requirement. +The exception to this is that only the destructor of an accessor of target +`access::target::local` destroys all containing elements and frees the memory. +Other accessor types do not own the memory, +so their destructor does not destroy the elements or free memory. + +Meeting the [ContiguousContainer](https://en.cppreference.com/w/cpp/named_req/ContiguousContainer) +named requirement means adding all the types, member functions, and operators +as defined by the [Container](https://en.cppreference.com/w/cpp/named_req/Container) +named requirement. +This includes, among other things, +adding a default constructor to all accessors. +Additionally, the iterator for the `accessor` container +is the pointer obtained by calling `accessor::get_pointer`. + +The change is not limited to 1D accessors. +For accessors of higher dimensions, +iterators are linearized according to standard SYCL linearization. + +When the access mode is `read`, +for a given underlying data type `T`, +the `ContiguousContainer` requirement is met for the type `const T`. + +## Examples + +```cpp +queue q; +static constexpr size_t r0 = 32; +static constexpr size_t r1 = 16; +const auto bufRange = range<2>{r0, r1}; +buffer buf{bufRange}; + +q.submit([&](handler& cgh) { + auto globalAcc = buf.get_access(cgh); + + auto localAcc = + accessor{ + range<1>{2 * r1}, cgh}; + + q.parallel_for(nd_range<2>{bufRange, range<2>{r1, 1}}, + [=](nd_item<2> i) { + using std::begin; + using std::end; + + const auto globalId = i.get_global_linear_id(); + const auto localId = i.get_local_linear_id(); + + // Initialize two consecutive values to the local ID + std::fill_n(begin(localAcc) + localId, 2, static_cast(localId)); + + i.barrier(); + if (localId == 0) { + // Accumulate local data and add global ID + localAcc[0] = std::accumulate(begin(localAcc), end(localAcc), globalId); + } + i.barrier(); + + // Iterator linearizes the 2D global buffer accessor + *(begin(globalAcc) + globalId) = localAcc[0]; + }); +}); + +{ + auto hostAcc = buf.get_access(); + for (auto val : hostAcc) { + std::cout << val << std::endl; + } +} +```