Support 0-dimensional buffers, accessors and kernels

[fknorr]

This PR implements support for 0-dimensional buffers, 0-dimensional global and local accessors as well as 0-dimensional kernels as proposed by the SYCL standard without relying on the feature being present in the SYCL implementation.

Most of the changed lines simply replace sycl::{range,id} with celerity::{range,id}. We need our own implementation here to support the 0-dimensional case independently from SYCL (in fact the new Celerity coordinate types can be arbitrary-dimensional, even though buffers and kernels are limited to Dims <= 3 as before).

0-dimensional accessors do not take range mappers on construction since every access must be an all-access anyway. To support the different signatures in CTAD the accessor type deduction logic was re-written from scratch.

The PR also gets rid of the broken backwards compatibility layer for sycl::item kernels.

Open Question

In SYCL, 0-dimensional accessors do not have an operator[](id<0>), but only operator*(). Should we support this (and also maybe range-mappers on 0-dimensional accessors) as an extension to ease generic programming on buffer dimensionality?

Edit: Yes we should.

Edit 2: Actually SYCL does not have pointer semantics but implicit conversions to reference-type for 0-dim accessors. We are currently undecided if we want to go the same route.

Edit 3: We are going to implement both, our sane API and also SYCL's for porting compatibility.

[psalz]

I did a first pass, looks good, thanks! Since you mentioned that you'll have to do the s/sycl::/celerity::/ replacement again anyway, could you maybe do it as a separate commit to make reviewing easier?

[github-actions]

clang-tidy made some suggestions

[psalz]

A few more notes!

[fknorr]

From live discussion with @PeterTh:

I didn't read the SYCL spec for 0-dimensional accessors properly. It defines a reference-conversion operator and an assignment operator instead of operator*:

  /* Available only when: (AccessMode != access_mode::atomic && Dimensions == 0) */
  operator reference() const;

  /* Available only when: (AccessMode != access_mode::atomic && Dimensions == 0) */
  const accessor& operator=(const value_type& other) const;

  /* Available only when: (AccessMode != access_mode::atomic && Dimensions == 0) */
  const accessor& operator=(value_type&& other) const;

This allows an accessor to be used like a T& reference in some contexts, but for a buffer of structs e.g. buffer<sycl::float4> would require a static cast before accessing a member, which is quite hideous. I like the pointer API (operator* and TODO operator->) much better in that regard.

We could now either copy what SYCL does or break with the standard, but IMO we should have some form of operator* or a get() function regardless for usability.

Regarding operator[] and range mappers for Dims == 0 we agree that we should have both. One minor weirdness would arise from this code being allowed since master-node host tasks are zero-dimensional, which is not apparent to the user otherwise:

buffer<int, 0> buf = 42;
q.submit([=](handler &cgh) {
    accessor acc(buf, cgh, one_to_one(), read_only);
    cgh.host_task(on_master_node, ...);
});

Edit: We're implementing both the SYCL API and our operator*/operator-> API.

[PeterTh]

I only looked at 22/44 files so far, but I'll probably not get to the rest soon so I'll post that part for now.

[github-actions]

clang-tidy made some suggestions

[PeterTh]

LGTM, but maybe run the GPUC2 benchmarks before merging and update the result files.