Skip to content

CP024 Default placeholders #122

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

Open
wants to merge 4 commits into
base: master
Choose a base branch
from
Open

CP024 Default placeholders #122

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
120b047
CP024 Default placeholders
ProGTX Mar 9, 2020
191abb8
MR Feedback
ProGTX Mar 13, 2020
fe8488c
Made constructors explicit
ProGTX Mar 31, 2020
104d6cb
Clarified `has_handler`
ProGTX Mar 31, 2020
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
3 changes: 2 additions & 1 deletion README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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_ |
| 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_ |
343 changes: 343 additions & 0 deletions default-placeholders/index.md
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,343 @@
# Default placeholders

| Proposal ID | CP024 |
|-------------|--------|
| Name | Default placeholders |
| Date of Creation | 9 March 2019 |
| Revision | 0.2 |
| Latest Update | 31 March 2020 |
| Target | SYCL Next (after 1.2.1) |
| Current Status | _Work in Progress_ |
| Reply-to | Peter Žužek <[email protected]> |
| Original author | Peter Žužek <[email protected]> |
| Contributors | Gordon Brown <[email protected]> |

## 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 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,
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 <typename dataT,
int dimensions,
access::mode accessMode,
access::target accessTarget = access::target::global_buffer,
access::placeholder isPlaceholder = access::placeholder::false_t>
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
Copy link
Contributor

Choose a reason for hiding this comment

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

Now that we allow an accessor to be null, we will need to define the behaviour when you pass a null accessor to require.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I've clarified that require throws cl::sycl::invalid_object_error if a null accessor is passed. Although since calling require is allowed on both placeholders and non-placeholders, it might be worth allowing it for null accessors as well, though that might introduce new subtle errors.

Copy link
Contributor

Choose a reason for hiding this comment

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

I think this is the right way to go, I think passing a null accessor to a kernel should be okay, as long as you don't dereference it, but calling require should fail as you're asking the runtime to create a data dependency on non-existent data. We could potentially allow it and have it be a no-op for generic programming sake, but if we have is_null I don't think that's necessary.

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 <typename dataT,
int dimensions,
access::mode accessMode,
access::target accessTarget,
access::placeholder isPlaceholder>
void require(accessor<dataT,
dimensions,
accessMode,
accessTarget,
isPlaceholder>
acc);
};
```

`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`.

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,
we propose deprecating it.

### Allow constructing host accessors from placeholders

Consider the following example:

```cpp
template<typename AccTypeA, typename AccTypeB>
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
Copy link
Contributor

@AerialMantis AerialMantis Mar 12, 2020 •
edited
Loading

Choose a reason for hiding this comment

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

We may also need to make a tweak to the buffer::get_access variant which takes no parameters, as that's currently defined to always return an accessor of access::target::host_buffer. If we want to be able to retrieve a placeholder accessor using this member function as well, we will need to disambiguate it. One option could be to require that the access::target be specified, another could be to have a different member function for this. Alternatively, we could just rely on the accessor constructor for this, as get_access is simply syntactic sugar over the constructor.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I'm not sure get_access should return a placeholder, it's probably better to always return a full accessor. In terms of verbosity, #100 would probably be sufficient for resolving that.

Copy link
Contributor

Choose a reason for hiding this comment

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

I think this is a good approach, this means that get_access is the full contract of associating the accessor with a buffer and handler, and if you want to create an accessor from a buffer without associating a handler with it you just use the constructor.

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.
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 <typename dataT,
int dimensions,
access::mode accessMode,
access::target accessTarget = access::target::global_buffer,
access::placeholder isPlaceholder = access::placeholder::false_t>
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)) &&
// (dimensions == 0)
explicit
accessor(buffer<dataT, 1> &bufferRef);

// 3
// Only available when: ((accessTarget == access::target::global_buffer) ||
// (accessTarget == access::target::constant_buffer)) &&
// (dimensions > 0)
accessor(buffer<dataT, dimensions> &bufferRef,
range<dimensions> accessRange,
id<dimensions> accessOffset = {});

// 4
// Only available when (accessTarget == access::target::host_buffer) &&
// ((otherTarget == access::target::global_buffer) ||
// (otherTarget == access::target::constant_buffer))
template <access::target otherTarget, access::placeholder otherPlaceholder>
explicit
accessor(accessor<dataT,
dimensions,
accessMode,
otherTarget,
otherPlaceholder>&
placeholderAcc);
};
```

### 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.
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 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.

```cpp
template <typename dataT,
int dimensions,
access::mode accessMode,
access::target accessTarget = access::target::global_buffer,
access::placeholder isPlaceholder = access::placeholder::false_t>
class accessor {
public:
...

// 1
bool is_null() const noexcept;

// 2
bool has_handler() const noexcept;

// 3
// Only available when ((accessTarget == access::target::global_buffer) ||
// (accessTarget == access::target::constant_buffer))
accessor<dataT,
dimensions,
accessMode,
access::target::host_buffer,
access::placeholder::false_t>
get_host_access() const;
};
```

## Examples

Simple vector addition:

```cpp
std::vector<int> a{1, 2, 3, 4, 5};
std::vector<int> b{6, 7, 8, 9, 10};

using read_acc =
accessor<int, 1, access::mode::read, access::target::global_buffer>;
using write_acc =
accessor<int, 1, access::mode::discard_write, access::target::global_buffer>;

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<int> bufA{bufRange};

accA = read_acc{bufA};
Copy link

@MikeDvorskiy MikeDvorskiy Mar 12, 2020 •
edited
Loading

Choose a reason for hiding this comment

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

  1. What's a practial sense of an accessor creation w/o buffer, if we re-create the same accessor later?
    M.b. it would make sense if we can "bind" a buffer with an "empty" accessor, which was created before?

  2. And more important suggestion: would be nice if we can get a buffer by an accessor which is associated
    with a command group.... or, better, "re-bind" that accessor to the host for getting result (in case if we have not a SYCL buffer in a scope where we want to get the result.)

Copy link
Contributor Author

Choose a reason for hiding this comment

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

  1. The main use case is for the accessor to be default constructible, so you could easily reuse the same data structures with different buffers. By assigning to it later you are effectively binding it. We could also have an accessor::bind member function, which would achieve the same thing as operator=. Or are you suggesting a buffer::bind member function?
  2. If I understand this correctly, you would capture a fully bound accessor once, store it on the host, and use the same accessor in multiple command groups and kernel invocations? I believe that's a separate, and larger, problem - if this proves to be a very useful use case we can look more into it. It seems a lot like USM, though. Also without USM, I don't think pointers are guaranteed to be persistent between kernel invocations, at least not in OpenCL.

Copy link

@MikeDvorskiy MikeDvorskiy Mar 13, 2020 •
edited
Loading

Choose a reason for hiding this comment

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

  1. Having an accessor::bind(buffer) member function is OK.
  2. I mean a situation when we have bound accessor with SYCL buffer once, we cannot get SYCL buffer back... What should I do for getting result on the host in the following case?
template<typename AccTypeA, typename AccTypeB>
void do_something_and_print_result(AccTypeA accA, AccTypeB accB)
{
   ....
    myQueue.submit([&](handler &cgh) {
       cgh.require(accA);
       cgh.require(accB);
       cgh.copy(accA, accB);
   });
   ...
   // For example, here we want a result access
   // But, in that scope we have not cl::sycl::buffer object (to get a host accessor), and there is 
   //not any way to get it.  Would be nice if there is some way to re-bind placholder accessor: 

/*   accB.rebind_to_host();  //here accB "becomes"  a host accessor  */

   for(auto i = 0; i < accB.get_count(); ++i)
       std::cout << accB[i] << " ";
   std::cout << std::endl;
   ...
}

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Sorry, I'm not able to access that issue.

That definitely sounds like a compelling use case. Modifying an accessor to change it from a device to a host one can be a bit problematic, especially since they're very different types. We briefly discussed that we could allow constructing a host accessor from a placeholder accessor. So your example would become - with this proposal and some of #100 - like this:

template<typename AccTypeA, typename AccTypeB>
void do_something_and_print_result(AccTypeA accA, AccTypeB accB)
{
   ....
    myQueue.submit([&](handler &cgh) {
       cgh.require(accA);
       cgh.require(accB);
       cgh.copy(accA, accB);
   });
   ...
   { // Remember to scope host accessors because they're essentially locks
       host_accessor hostAcc{accB};
       // or: auto hostAcc = accB.get_host_access();
       for(auto i = 0; i < hostAcc.get_count(); ++i)
           std::cout << hostAcc[i] << " ";
       std::cout << std::endl;
       ...
   }
}

We do not have clear consensus on whether this should be permitted or not, but initial investigation reveals this might be OK.

Copy link

Choose a reason for hiding this comment

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

  1. Sounds great! "Constructing a host accessor from a placeholder accessor" should be enough for our use cases... Thank you.
    Does some formal step require from me to it will be adopted in the spec? Or that example is enough?
  2. I'll try to clarify how you can get an access to https://gitlab.devtools.intel.com/SYCL/specbugs/issues/65

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I've added your example to the proposal, if that's OK, and added a host accessor constructor from a placeholder, and also a get_host_access function to placeholders.

For spec adoption, proposals are usually brought to the Khronos group, either on GitHub or internally, and your support at that point helps a lot.

Copy link

Choose a reason for hiding this comment

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

Good! Thank you.

Copy link

@MikeDvorskiy MikeDvorskiy Mar 30, 2020 •
edited
Loading

Choose a reason for hiding this comment

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

I faced to another issue connected with usage of a placeholder accessor - there is not a way(API) to change access mode for given accessor. Please have a look following example:

template<typename AccTypeA, typename AccTypeB>
void do_something_and_print_result(AccTypeA accA, AccTypeB accB)
{
   ....
    myQueue.submit([&](handler &cgh) {
       cgh.require(accA);
       cgh.require(accB);
       cgh.copy(accA, accB);
   });
   ...
   // For example, here we want to do some computation  for data "B", a kind of reduction, for example. We need just 
   //"Read" access for "B", but passed accB has "Write" access mode. So, would be great to have possibility to create (or change) a new accessor with necessary access mode, like this:

read_accessor accB_read{accB}; 

   ...
}

Copy link
Contributor Author

Choose a reason for hiding this comment

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

#100 talks about allowing accessors to be implicitly converted to accessor<const dataT> and also to a read-only accessor (because read-only should be semantically the same as const dataT).

I'm not sure how general this should be, because even just going in the other direction gets very tricky very fast, you get into problems with multiple accessor resolution.

Copy link

@MikeDvorskiy MikeDvorskiy Apr 7, 2020 •
edited
Loading

Choose a reason for hiding this comment

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

Thanks! It seems that enough....

More one important use-case which I faced to:

template<typename T /* this is some accessor */ >
void foo(T& in)
{
   // For example, here we want to do some computation  for part of data - we want to skip first N elements. So, we 
   // would like to set "offset = N", but there is no API for that(!)
   
   in.set_offset(N); 
   //or accessor new_acc(in, N);

   myQueue.submit([&](handler &cgh) {
     ...
   }
   ...

}

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<int> 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
buffer<int> bufC{bufRange};
accC = read_acc{bufC};
myQueue.submit([&](handler &cgh) {
cgh.require(accA);
cgh.require(accB);
cgh.require(accC);
cgh.parallel_for<class vec_add>(bufRange,
[=](id<1> i) { accC[i] = accA[i] + accB[i]; });
});
```