[SYCL] Cannot use derived accessor objects

[Ruyk]

The following code:

#include <array>

#include <CL/sycl.hpp>

constexpr cl::sycl::access::mode sycl_read = cl::sycl::access::mode::read;
constexpr cl::sycl::access::mode sycl_write = cl::sycl::access::mode::write;
constexpr size_t N = 10u;

template<typename T>
class SimpleVadd;

template<typename T>
using my_nice_alias = cl::sycl::accessor<T, 1, sycl_write, 
                                        cl::sycl::access::target::global_buffer>;

template<typename T>
struct my_accessor : public my_nice_alias<T> {

  using my_nice_alias<T>::my_nice_alias;
};

int main() {
        using T = int;
  using array_type = std::array<T, N>;

        array_type VA, VB, VC;

  for (size_t i = 0; i < N; i++) {
    VA[i] = VB[i] = i;
  }

  {
                cl::sycl::queue deviceQueue;
                cl::sycl::range<1> numOfItems{N};
                cl::sycl::buffer<T, 1> bufferA(VA.data(), numOfItems);
                cl::sycl::buffer<T, 1> bufferB(VB.data(), numOfItems);
                cl::sycl::buffer<T, 1> bufferC(VC.data(), numOfItems);

                deviceQueue.submit([&](cl::sycl::handler& cgh) {
                        auto accessorA = bufferA.template get_access<sycl_read>(cgh);
                        auto accessorB = bufferB.template get_access<sycl_read>(cgh);
#if NORMAL_ACCESSOR
                        auto accessorC = bufferC.template get_access<sycl_write>(cgh);
#else
                        my_accessor<int> accessorC(bufferC, cgh);
#endif

                        cgh.parallel_for<class SimpleVadd<T>>(numOfItems,
                        [=](cl::sycl::id<1> wiID) {
                                accessorC[wiID] = accessorA[wiID] + accessorB[wiID];
                        });
                });
        }
        return !(VC[0] == (VA[0] + VB[0]));
}

produces incorrrect results when NORMAL_ACCESSOR is not defined, but works fine when it is.

The main difference between the two code paths is that we use a derived class from the SYCL accessor, which is valid from the SYCL specification.
The generated SPIR-V binary doesn't recognize the derived class (my_accessor) as an accessor, and fails to generate the correct code (treats my_accessor as a normal struct). However, no error is reported: It just generates incorrect results at runtime.

Seems this could be to the fact that

llvm/clang/lib/Sema/SemaSYCL.cpp

Line 1087 in bae2b75

bool Util::isSyclAccessorType(const QualType &Ty) {

checks the name of the type to see if it matches the name accessor, which seems to me its a bit fragile for long-term maintenance?
Libraries or users can derive from accessors to create custom types.
I suspect a similar problem may occur if users derive from the stream class.