diff --git a/placeholder_accessors/index.md b/placeholder_accessors/index.md index ed8fe92..2289250 100644 --- a/placeholder_accessors/index.md +++ b/placeholder_accessors/index.md @@ -9,7 +9,7 @@ | Reply-to | Ruyman Reyes | | Original author | Ruyman Reyes | | Requirements | CP001 | -| Contributors | Gordon Brown , Victor Lomuller , Mehdi Goli , Peter Zuzek , Luke Iwanski | +| Contributors | Gordon Brown , Victor Lomuller , Mehdi Goli , Peter Žužek , Luke Iwanski , Michael Haidl | ## Overview @@ -172,7 +172,8 @@ except for the aforementioned modifications. ### When `is_placeholder` returns true -The accessor API features constructors that don't take the handler parameter. +The accessor API features constructors that don't take the handler parameter +and/or memory object as a constructor. In addition, a new method to obtain a normal accessor from the placeholder accessor is provided. @@ -192,6 +193,51 @@ The handler gains a new method, that registers the requirements of the placeholder accessor on the given command group. +### Placeholder `accessor` without a buffer + +Not providing either the handler or the memory object +makes the accessor default constructible, +creating a "null" accessor. +A null accessor can be bound to a buffer +by constructing a new placeholder accessor with the buffer +and assigning it to the null accessor. +Example: + +```cpp +accessor acc0; + +acc0 = accessor(buf); +``` + +To simplify placeholder accessor construction, +a new free function `cl::sycl::codeplay::make_placeholder_accessor` is provided, +which takes a buffer as an argument: + +```cpp +template +accessor +make_placeholder_accessor(buffer& bufferRef); +``` + +A null accessor does not need to be registered with the command group, +i.e. there is no need to call `require()` on it. +It is allowed to pass a null accessor to a kernel, +as long as it's not dereferenced inside the kernel, +otherwise an out-of-bounds access will occur, +likely leading to a segmentation fault or similar. +If the null accessor has been bound to a buffer, +it is no longer a null accessor, +which means it has to be registered with the command group +and can be dereferenced inside the kernel. + +An accessor can be checked for the existence of an associated buffer +by calling the member function `is_null`. + +|Member function |Description | +|---------------------|---------------------------------------------------------| +|bool is_null() const |Returns false if the accessor is associated with a buffer, true otherwise.| + [1]: https://github.com/codeplaysoftware/sycl-blas "SYCL-BLAS" [2]: https://github.com/lukeiwanski/tensorflow "TensorFlow/Eigen"