Is it allowed to wrap accessors?

[psalz]

The SYCL 1.2.1 Rev 6 spec states that all user defined data structures that are passed as arguments to kernels must have C++11 standard layout (section 3.10). Intel is proposing to relax this to trivially copyable. However, I cannot find any clarification on whether the built-in argument types, i.e., stream, sampler, and in particular accessor, must also fulfill this requirement, and by extension, whether it is allowed to wrap these types inside user defined data structures.

There is another sentence in section 3.10 that might imply it not being allowed (emphasis is mine):

The only way of passing pointers to a kernel is through the cl::sycl::accessor class, which supports the cl::sycl::buffer and cl::sycl::image classes. No hierarchical structures of these classes are supported and any other data containers need to be converted to the SYCL data management classes using the SYCL interface.

However from my reading I think this refers to buffer and image, not accessor.

In any case, out of the four SYCL implementations I tested, none provides accessors that are both standard layout and trivially copyable:

Impl	is_standard_layout	is_trivially_copyable
hipSYCL	false	false
ComputeCpp	true	false
Intel	false	false
triSYCL	true	false

Since these requirements are defined recursively, any user defined type wrapping an accessor from one of these implementations may thus not have standard layout (and certainly won't be trivially copyable).

I would actually find it quite puzzling if wrapping accessors were not allowed, as it seems a rather useful pattern to me. In particular since as far as I understand, wrapping and later "hydrating" accessors was one of the main justifications for introducing placeholder accessors in the first place. This pattern is currently e.g. extensively used in Codeplay's SYCL-BLAS.

Can you clarify this?

[keryell]

I guess the real problem that bothers you with the accessors is not whether it can be passed to a kernel as any other class, but if at the end it is understood as an accessor with the expected semantics. My understanding it that ComputeCpp and DPC++ allow classes to inherit from an accessor and to behave like an accessor. It is not the case in triSYCL with the legacy outliner where the type of the class has to exactly match an accessor. For hipSYCL I do not know.

[psalz]

I guess the real problem that bothers you with the accessors is not whether it can be passed to a kernel as any other class, but if at the end it is understood as an accessor with the expected semantics.

Well, that's how I'd like to use it, yes. But what I'm asking here is what the spec's position on it is, and whether it can give me any guarantees. I'd rather not build my own APIs using patterns that will be declared UB later down the line...

[keryell]

After thinking about it in some other context (in the private https://gitlab.khronos.org/sycl/Specification/merge_requests/425 for my Khronos friends), I think what you are asking for is required just to fit in some modern metaprogramming: having a generic SYCL executor where you pack all the accessors in a std::tuple and you use them from inside the kernel. Since it is already in 2 implementations, at least we know that it is not science-fiction. :-)

[keryell]

While creating an internal issue for this, I realized that there is already the internal https://gitlab.khronos.org/sycl/Specification/issues/225

[keryell]

Even worse I just realize I have made already presentations at SuperComputing 2017 showing why we need this to implement generic executors https://www.khronos.org/developers/library/sc17 https://www.khronos.org/assets/uploads/developers/library/2017-supercomputing/Xilinx-triSYCL-complete_Nov17.pdf slides 47—49. Alzheimer is my friend... :-(

[illuhad]

My understanding it that ComputeCpp and DPC++ allow classes to inherit from an accessor and to behave like an accessor. It is not the case in triSYCL with the legacy outliner where the type of the class has to exactly match an accessor. For hipSYCL I do not know.

hipSYCL does not use any compiler magic to implement accessors, so I would expect either wrapping them in another class or using them as base class to work.

[keryell]

I think we can close this since it has been clarified with https://github.com/KhronosGroup/SYCL-Docs/pull/173 in SYCL 2020 https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:kernel.parameter.passing 4.12.4. Rules for parameter passing to kernels as allowed to inherit from an accessor according to item:

• A class type S with a non-virtual base class of type T is a legal parameter type if T is a legal parameter type and if S would otherwise be a legal parameter type aside from this base class.