Extension to make accessors optional in SYCL

[keryell]

Accessors are very useful to precise the use of some resources, such as reading a buffer, accessing a pipe in write blocking mode, passing some local memory, etc.

Accessors allow the unique SYCL feature of defining implicitly the dependencies of an asynchronous task graph, without requiring specific compiler support for this.

In some cases, a programmer may not need it:

• porting some code from CUDA or from other heterogeneous frameworks where the code has already the dependency and data management;
• not wanting this implicit dependency management and wanting to use C++ continuation-style programming, with .then() and so on or using in the context of a C++ SG1 executor;
• a system has Fine-Grain System SVM, such as a common memory between host and device.

So it could be possible to access buffers or pipes directly from kernels or host. But there might not be some sanity check at runtime.

Or we could have a more compiler-oriented approach as suggested by @MathiasMagnus where the compiler can figure out the dependencies and synthesize the equivalent of the missing accessors. Which means that the code would not work on CPU without a specific compiler to do a job similar to some automatic parallelizers. Such compilers will not being able to make a difference between some read_write, write and discard_write modes for example and lead to less optimal code.

[MathiasMagnus]

I have already posted this question on the KhronosDev Slack channel, but nobody provided feedback. Do you happen to know of a .then() implementation to cl::Event or cl::sycl::event that accepts arbitrary C++ callables and feeds it to the underlying OpenCL runtime properly? (The question is a little more fleshed out on Slack)

What EXACTLY is the semantic difference between write and discard_write access modes? Do I understand correctly that the only difference is the promise from the programmer that the entire (sub-)buffer is written to or not? If I overwrite every element, then it's discard_write, but if I do the same and leave some elements untouched it will become garbage, right? I believe only this latter two is what is less trivial to infer by the compiler, and even this isn't impossible to figure out, at least in the simpler (majority?) of cases.

[keryell]

We have some ideas about providing lower-level continuation interface. I do not know whether there is an equivalent implemented already in the OpenCL realm.

In 4.7.6.2 Access modes:

If a command group contains only discard write mode accesses to a buffer, then the previous contents of the buffer (or sub-range of the buffer, if provided) are not preserved. If a user wants to modify only certain parts of a buffer, preserving other parts of the buffer, then the user should specify the exact sub-range of modification of the buffer.

Yes, if you forget to initialize some values you will have undefined behaviour.

PS: I have just registered to the Khronos Slack. :-)

[keryell]

This is related to the new SYCL 2020 USM.