CP024 Default placeholders

[ProGTX]

1. Deprecate access::placeholder
2. Extend placeholder functionality to ignore template parameter
   • Still applies only to global and constant buffers
3. Updated handler::require
4. New accessor constructors
   • Default constructor
   • From a buffer, 0-dim
   • From a buffer, ranged
5. New access member functions
   • is_null
   • has_handler
6. Vector addition example

[keryell]

That looks good.

A minor point I see is that before by default an accessor constructor without a CGH was a host accessor. Now it is no longer the case, since it might be a default placeholder accessor to be used in a kernel instead.

[MikeDvorskiy]

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

[ProGTX]

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.

[MikeDvorskiy]

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;
   ...
}

[ProGTX]

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.

[MikeDvorskiy]

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

[ProGTX]

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.

[MikeDvorskiy]

Good! Thank you.

[MikeDvorskiy]

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}; 

   ...
}

[ProGTX]

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

[MikeDvorskiy]

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) {
     ...
   }
   ...

}

[AerialMantis]

Looks great, thanks! Just a few suggestions for the wording.

[AerialMantis]

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.

[ProGTX]

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.

[AerialMantis]

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.

[AerialMantis]

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

[ProGTX]

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.

[AerialMantis]

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.

[ProGTX]

A minor point I see is that before by default an accessor constructor without a CGH was a host accessor. Now it is no longer the case, since it might be a default placeholder accessor to be used in a kernel instead.

@keryell You're right that the default when constructing an accessor will now be a placeholder instead of a host accessor, but this is only really important for CTAD, otherwise you'd still need to specify target::host_buffer. Allowing CTAD to determine either a device accessor or a host accessor based on whether the handler was provided seems dangerous to us - it is already a common mistake to call get_access without a handler, which gives you a host accessor in a command group, which is UB. At least with a placeholder you still have the option of calling require (which is valid for placeholders and non-placeholders according to this proposal), which gets rid of the UB.

There is even some interest in making host accessors a separate type, as an extension of the work here and in #100.

[keryell]

Allowing CTAD to determine either a device accessor or a host accessor based on whether the handler was provided seems dangerous to us - it is already a common mistake to call get_access without a handler, which gives you a host accessor in a command group, which is UB.

The compiler should warn about it at least.
But this is what I do in all the example I show because it is shorter...

[ProGTX]

The compiler should warn about it at least.
But this is what I do in all the example I show because it is shorter...

I think that calling buf.get_access<mode>() is still fine, I'm just saying that sometimes this happens because the user forgets to pass the handler - they wanted a device accessor, but got a host one instead, so with new proposals I'm keen on making more of a distinction between host and device accessors. I'm not sure if the compiler has enough information to warn on this.