intel / llvm

Intel staging area for llvm.org contribution. Home for Intel LLVM-based projects.
Other
1.25k stars 738 forks source link

accessor should export its template arguments / template matching error on accessor<T,int,auto...> #10057

Open stdale-intel opened 1 year ago

stdale-intel commented 1 year ago

When storing SYCL accessors inside a struct, it is important to be able to access the accessor's template parameters. (Refs: https://github.com/intel/llvm/issues/9434, https://github.com/KhronosGroup/SYCL-Docs/issues/421).

The accessor documentation shows a definition:

template <typename DataT, int Dimensions = 1,
          access_mode AccessMode =
              (std::is_const_v<DataT> ? access_mode::read
                                      : access_mode::read_write),
          target AccessTarget = target::device,
          access::placeholder isPlaceholder = access::placeholder::false_t>
class accessor

However, this definition is defective, because the accessor does not re-export its access mode or target.

Further, attempting to work around this limitation exposes an implementation flaw in DPCPP. The following code works generally and also succeeds with hipSYCL. However, it fails to compile with dpcpp:

#include <sycl/sycl.hpp>

template<typename, typename>
struct apply_accessor {};

template<template<typename,auto...> typename Acc,
         typename T1, typename T, auto... A>
struct apply_accessor<Acc<T1, A...>, T> {
    using type = Acc<T, A...>;
};

struct A {
    using X = sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::device>;//, sycl::access::placeholder::false_t>;

    using Y = typename apply_accessor<X, float>::type;
};

int main() {
    A a;
}

Copy of filing from: https://github.com/KhronosGroup/SYCL-Docs/issues/433

0x12CC commented 4 months ago

@frobnitzem, our accessor class template doesn't match your specialization since it differs from the one in AdaptiveCpp. I think it's due to our optional sycl::ext::oneapi::accessor_property_list<> parameter. As a workaround, you can add one more specialization:

#include <sycl/sycl.hpp>

template <typename, typename> struct apply_accessor {};

template <template <typename, auto...> typename Acc, typename T1, typename T,
          auto... A>
struct apply_accessor<Acc<T1, A...>, T> {
  using type = Acc<T, A...>;
};

template <template <typename, auto A, auto B, auto C, auto D, typename E>
          typename Acc,
          typename T1, typename T, auto A, auto B, auto C, auto D, typename E>
struct apply_accessor<Acc<T1, A, B, C, D, E>, T> {
  using type = Acc<T, A, B, C, D, E>;
};

struct A {
  using X = sycl::accessor<
      int, 1, sycl::access::mode::read_write,
      sycl::access::target::device>; //, sycl::access::placeholder::false_t>;

  using Y = typename apply_accessor<X, float>::type;
};

int main() { A a; }

@Pennycook, could you please take a look at this? It looks like the example code fails to compile due to the sycl_ext_oneapi_accessor_properties extension but I'm not sure if we can fix this.

Pennycook commented 3 months ago

I agree that accessor should expose its template arguments as members, and that's something we should try to address via Khronos.

I also agree that our accessor properties extension appears to be the root cause here, and that it's going to be difficult to fix. Ultimately, we're going to need to find a way to deprecate this extension and bring it in line with our rework of the compile-time properties extension.