Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Default constructed accessors #89

Closed
wants to merge 2 commits into from
Closed

Default constructed accessors #89

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
7bb277e
Default constructed accessors
ProGTX Jan 24, 2019
b8a9189
Minor fixes
ProGTX Jan 25, 2019
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
50 changes: 48 additions & 2 deletions placeholder_accessors/index.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@
| Reply-to | Ruyman Reyes <[email protected]> |
| Original author | Ruyman Reyes <[email protected]> |
| Requirements | CP001 |
| Contributors | Gordon Brown <[email protected]>, Victor Lomuller <[email protected]>, Mehdi Goli <[email protected]>, Peter Zuzek <[email protected]>, Luke Iwanski <[email protected]> |
| Contributors | Gordon Brown <[email protected]>, Victor Lomuller <[email protected]>, Mehdi Goli <[email protected]>, Peter Žužek <[email protected]>, Luke Iwanski <[email protected]>, Michael Haidl <[email protected]> |

## Overview

Expand Down Expand Up @@ -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.
Expand All @@ -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.
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

a "null" accessor -> a "null" or "unbound" placeholder accessor. This represents just the placeholder (i.e. the space that will be occupied by the data type) that needs to be associated with a buffer at command group submission time.

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<T, dim, mode, target, access::placeholer::true_t> acc0;

acc0 = accessor<T, dim, mode, target, access::placeholer::true_t>(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 <access::mode mode, access::target target, typename T, int dim,
typename AllocatorT>
accessor<T, dim, mode, target, access::placeholder::true_t>
make_placeholder_accessor(buffer<T, dim, AllocatorT>& 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.
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is quite ambiguous. Is better simply to say a device-specific asynchronous error

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.
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is also a bit confusing. A null accessor bound to a buffer, is then bound to that buffer forever? can it re-assigned? can it be assigned back to null? What happens if two command groups submit a different require bounding the same null accessor to two separate buffers?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The main point of this MR is that the require function doesn't do the binding - it doesn't modify the accessor at all. So a null accessor can be "bound" by assigning to it a placeholder that already has an associated buffer.


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.|
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What is the expected output of is_null in the following cases?

Case 1:

accessor<T, dim, mode, target, access::placeholer::true_t> acc0;

q.submit([&](handler& h) {
      h.require(acc0, bufA);
});

// What is the output here?
std::cout << acc.is_null() << std::endl;

Case 2:

accessor<T, dim, mode, target, access::placeholer::true_t> acc0;

std::thread t([](){
  q.submit([&](handler& h) {
      h.require(acc0, bufA);
  });
});
// What is the output here?
std::cout << acc.is_null() << std::endl;
t.join();

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This PR doesn't have that overload of the require function. The existing function doesn't modify the state of the accessor, and having an overload that would modify it is a dangerous inconsistency. That also means that both your cases would print true.

h.require(acc0, bufA);
// is instead
acc0 = accessor<T, dim, mode, target, access::placeholer::true_t>(bufA);
h.require(acc0);
// or
acc0 = make_placeholder_accessor<mode, target>(bufA);
h.require(acc0);



[1]: https://github.com/codeplaysoftware/sycl-blas "SYCL-BLAS"
[2]: https://github.com/lukeiwanski/tensorflow "TensorFlow/Eigen"
Expand Down