[SYCL] Fix several errors in accessor

This patch fixes 2 constructors, fixes get_count() and get_size() methods,
adds operator== and operator!=, fixes hash implementation,
adds static_asssert to range and id classes to have early check for wrong
dimension, fixes few other errors.

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

Author: v-klochkov

sycl/include/CL/sycl/accessor2.hpp

@@ -212,7 +212,7 @@ class accessor_common {
       MIDs[0] = Index;
     }
 
-    template <int CurDims = SubDims, typename = enable_if_t<CurDims != 1>>
+    template <int CurDims = SubDims, typename = enable_if_t<(CurDims > 1)>>
     AccessorSubscript<CurDims - 1> operator[](size_t Index) {
       MIDs[Dims - CurDims] = Index;
       return AccessorSubscript<CurDims - 1>(MAccessor, MIDs);
@@ -282,19 +282,19 @@ class accessor :
 
     size_t Result = 0;
     for (int I = 0; I < Dims; ++I)
-      Result = Result * getOrigRange()[I] + getOffset()[I] + Id[I];
+      Result = Result * getMemoryRange()[I] + getOffset()[I] + Id[I];
     return Result;
   }
 
 #ifdef __SYCL_DEVICE_ONLY__
 
   id<AdjustedDim> &getOffset() { return impl.Offset; }
-  range<AdjustedDim> &getRange() { return impl.AccessRange; }
-  range<AdjustedDim> &getOrigRange() { return impl.MemRange; }
+  range<AdjustedDim> &getAccessRange() { return impl.AccessRange; }
+  range<AdjustedDim> &getMemoryRange() { return impl.MemRange; }
 
   const id<AdjustedDim> &getOffset() const { return impl.Offset; }
-  const range<AdjustedDim> &getRange() const { return impl.AccessRange; }
-  const range<AdjustedDim> &getOrigRange() const { return impl.MemRange; }
+  const range<AdjustedDim> &getAccessRange() const { return impl.AccessRange; }
+  const range<AdjustedDim> &getMemoryRange() const { return impl.MemRange; }
 
   detail::AccessorImplDevice<AdjustedDim> impl;
 
@@ -305,8 +305,8 @@ class accessor :
     MData = Ptr;
     for (int I = 0; I < AdjustedDim; ++I) {
       getOffset()[I] = Offset[I];
-      getRange()[I] = AccessRange[I];
-      getOrigRange()[I] = MemRange[I];
+      getAccessRange()[I] = AccessRange[I];
+      getMemoryRange()[I] = MemRange[I];
     }
     // In case of 1D buffer, adjust pointer during initialization rather
     // then each time in operator[] or get_pointer functions.
@@ -317,9 +317,9 @@ class accessor :
   PtrType getQualifiedPtr() const { return MData; }
 #else
 
-  using AccessorBaseHost::getRange;
+  using AccessorBaseHost::getAccessRange;
   using AccessorBaseHost::getOffset;
-  using AccessorBaseHost::getOrigRange;
+  using AccessorBaseHost::getMemoryRange;
 
   char padding[sizeof(detail::AccessorImplDevice<AdjustedDim>) +
                sizeof(PtrType) - sizeof(detail::AccessorBaseHost)];
@@ -332,17 +332,15 @@ class accessor :
 
 public:
   using value_type = DataT;
-  using reference = typename detail::PtrValueType<DataT, AS>::type &;
+  using reference = DataT &;
   using const_reference = const reference;
 
   template <int Dims = Dimensions>
-  accessor(buffer<DataT, 1> &BufferRef,
-           enable_if_t<((!IsPlaceH && IsHostBuf) ||
+  accessor(enable_if_t<((!IsPlaceH && IsHostBuf) ||
                         (IsPlaceH && (IsGlobalBuf || IsConstantBuf))) &&
-                       Dims == 0>)
+                       Dims == 0, buffer<DataT, 1> > &BufferRef)
 #ifdef __SYCL_DEVICE_ONLY__
-      : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.MemRange)
-      {
+      : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.MemRange) {
 #else
       : AccessorBaseHost(
             /*Offset=*/{0, 0, 0},
@@ -351,15 +349,17 @@ class accessor :
             detail::getSyclObjImpl(BufferRef).get(), AdjustedDim,
             sizeof(DataT)) {
     detail::EventImplPtr Event =
-        detail::Scheduler::getInstance().addHostAccessor(this);
+        detail::Scheduler::getInstance().addHostAccessor(
+            AccessorBaseHost::impl.get());
     Event->wait(Event);
 #endif
   }
 
   template <int Dims = Dimensions>
   accessor(
-      buffer<DataT, 1> &BufferRef, handler &CommandGroupHandler,
-      enable_if_t<(!IsPlaceH && (IsGlobalBuf || IsConstantBuf)) && Dims == 0>)
+      buffer<DataT, 1> &BufferRef,
+      enable_if_t<(!IsPlaceH && (IsGlobalBuf || IsConstantBuf)) && Dims == 0,
+                   handler> &CommandGroupHandler)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.MemRange) {
   }
@@ -458,13 +458,13 @@ class accessor :
 
   constexpr bool is_placeholder() const { return IsPlaceH; }
 
-  size_t get_size() const { return getRange().size() * sizeof(DataT); }
+  size_t get_size() const { return getMemoryRange().size() * sizeof(DataT); }
 
-  size_t get_count() const { return getRange().size(); }
+  size_t get_count() const { return getMemoryRange().size(); }
 
   template <int Dims = Dimensions, typename = enable_if_t<(Dims > 0)>>
   range<Dimensions> get_range() const {
-    return detail::convertToArrayOfN<Dimensions, 1>(getRange());
+    return detail::convertToArrayOfN<Dimensions, 1>(getAccessRange());
   }
 
   template <int Dims = Dimensions, typename = enable_if_t<(Dims > 0)>>
@@ -475,7 +475,7 @@ class accessor :
   template <int Dims = Dimensions,
             typename = enable_if_t<IsAccessAnyWrite && Dims == 0>>
   operator RefType() const {
-    const size_t LinearIndex = getLinearIndex(id<Dimensions>());
+    const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return *(getQualifiedPtr() + LinearIndex);
   }
 
@@ -570,6 +570,9 @@ class accessor :
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return constant_ptr<DataT>(getQualifiedPtr() + LinearIndex);
   }
+
+  bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
+  bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
 };
 
 // Local accessor
@@ -604,8 +607,8 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
 #ifdef __SYCL_DEVICE_ONLY__
   detail::LocalAccessorBaseDevice<AdjustedDim> impl;
 
-  sycl::range<AdjustedDim> &getSize() { return impl.AccessRange; }
-  const sycl::range<AdjustedDim> &getSize() const { return impl.AccessRange; }
+  sycl::range<AdjustedDim> &getSize() { return impl.MemRange; }
+  const sycl::range<AdjustedDim> &getSize() const { return impl.MemRange; }
 
   void __init(PtrType Ptr, range<AdjustedDim> AccessRange,
               range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
@@ -719,6 +722,9 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
   local_ptr<DataT> get_pointer() const {
     return local_ptr<DataT>(getQualifiedPtr());
   }
+
+  bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
+  bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
 };
 
 // Image accessor
@@ -815,9 +821,10 @@ struct hash<cl::sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
     // Hash is not supported on DEVICE. Just return 0 here.
     return 0;
 #else
-    std::shared_ptr<cl::sycl::detail::AccessorBaseHost> AccBaseImplPtr =
-        cl::sycl::detail::getSyclObjImpl(A);
-    return hash<decltype(AccBaseImplPtr)>()(AccBaseImplPtr);
+    // getSyclObjImpl() here returns a pointer to either AccessorImplHost
+    // or LocalAccessorImplHost depending on the AccessTarget.
+    auto AccImplPtr = cl::sycl::detail::getSyclObjImpl(A);
+    return hash<decltype(AccImplPtr)>()(AccImplPtr);
 #endif
   }
 };

sycl/include/CL/sycl/detail/accessor_impl.hpp

@@ -27,12 +27,19 @@ namespace detail {
 
 template <int Dims> class AccessorImplDevice {
 public:
-  AccessorImplDevice(id<Dims> Offset, range<Dims> Range, range<Dims> OrigRange)
-      : Offset(Offset), AccessRange(Range), MemRange(OrigRange) {}
+  AccessorImplDevice(id<Dims> Offset, range<Dims> AccessRange,
+                     range<Dims> MemoryRange)
+      : Offset(Offset), AccessRange(AccessRange), MemRange(MemoryRange) {}
 
   id<Dims> Offset;
   range<Dims> AccessRange;
   range<Dims> MemRange;
+
+  bool operator==(const AccessorImplDevice &Rhs) const {
+    return (Offset == Rhs.Offset &&
+            AccessRange == Rhs.AccessRange &&
+            MemRange == Rhs.MemRange);
+  }
 };
 
 template <int Dims> class LocalAccessorBaseDevice {
@@ -43,14 +50,18 @@ template <int Dims> class LocalAccessorBaseDevice {
   range<Dims> AccessRange;
   range<Dims> MemRange;
   id<Dims> Offset;
+
+  bool operator==(const LocalAccessorBaseDevice &Rhs) const {
+    return (AccessRange == Rhs.AccessRange);
+  }
 };
 
 class AccessorImplHost {
 public:
-  AccessorImplHost(id<3> Offset, range<3> Range, range<3> OrigRange,
+  AccessorImplHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
                    access::mode AccessMode, detail::SYCLMemObjT *SYCLMemObject,
                    int Dims, int ElemSize)
-      : MOffset(Offset), MRange(Range), MOrigRange(OrigRange),
+      : MOffset(Offset), MAccessRange(AccessRange), MMemoryRange(MemoryRange),
         MAccessMode(AccessMode), MSYCLMemObj(SYCLMemObject), MDims(Dims),
         MElemSize(ElemSize) {}
 
@@ -59,16 +70,16 @@ class AccessorImplHost {
       BlockingEvent->setComplete();
   }
   AccessorImplHost(const AccessorImplHost &Other)
-      : MOffset(Other.MOffset), MRange(Other.MRange),
-        MOrigRange(Other.MOrigRange), MAccessMode(Other.MAccessMode),
+      : MOffset(Other.MOffset), MAccessRange(Other.MAccessRange),
+        MMemoryRange(Other.MMemoryRange), MAccessMode(Other.MAccessMode),
         MSYCLMemObj(Other.MSYCLMemObj), MDims(Other.MDims),
         MElemSize(Other.MElemSize) {}
 
   id<3> MOffset;
   // The size of accessing region.
-  range<3> MRange;
+  range<3> MAccessRange;
   // The size of memory object this requirement is created for.
-  range<3> MOrigRange;
+  range<3> MMemoryRange;
   access::mode MAccessMode;
 
   detail::SYCLMemObjT *MSYCLMemObj;
@@ -85,22 +96,23 @@ using AccessorImplPtr = std::shared_ptr<AccessorImplHost>;
 
 class AccessorBaseHost {
 public:
-  AccessorBaseHost(id<3> Offset, range<3> Range, range<3> OrigRange,
+  AccessorBaseHost(id<3> Offset, range<3> AccessRange, range<3> MemoryRange,
                    access::mode AccessMode, detail::SYCLMemObjT *SYCLMemObject,
                    int Dims, int ElemSize) {
-    impl = std::make_shared<AccessorImplHost>(
-        Offset, Range, OrigRange, AccessMode, SYCLMemObject, Dims, ElemSize);
+    impl = std::make_shared<AccessorImplHost>(Offset, AccessRange, MemoryRange,
+                                              AccessMode, SYCLMemObject,
+                                              Dims, ElemSize);
   }
 
 protected:
   id<3> &getOffset() { return impl->MOffset; }
-  range<3> &getRange() { return impl->MRange; }
-  range<3> &getOrigRange() { return impl->MOrigRange; }
+  range<3> &getAccessRange() { return impl->MAccessRange; }
+  range<3> &getMemoryRange() { return impl->MMemoryRange; }
   void *getPtr() { return impl->MData; }
 
   const id<3> &getOffset() const { return impl->MOffset; }
-  const range<3> &getRange() const { return impl->MRange; }
-  const range<3> &getOrigRange() const { return impl->MOrigRange; }
+  const range<3> &getAccessRange() const { return impl->MAccessRange; }
+  const range<3> &getMemoryRange() const { return impl->MMemoryRange; }
   void *getPtr() const { return const_cast<void *>(impl->MData); }
 
   template <class Obj>
@@ -135,7 +147,10 @@ class LocalAccessorBaseHost {
 
   int getNumOfDims() { return impl->MDims; }
   int getElementSize() { return impl->MElemSize; }
+
 protected:
+  template <class Obj>
+  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
 
   std::shared_ptr<LocalAccessorImplHost> impl;
 };

sycl/include/CL/sycl/detail/array.hpp

@@ -19,6 +19,7 @@ template <int dimensions> class range;
 namespace detail {
 
 template <int dimensions = 1> class array {
+  static_assert(dimensions >= 1, "Array cannot be 0-dimensional.");
 public:
   array() : common_array{0} {}
 

sycl/include/CL/sycl/handler2.hpp

@@ -234,11 +234,11 @@ class handler {
 
           MArgs.emplace_back(
               detail::ArgDesc(detail::kernel_param_kind_t::kind_std_layout,
-                              &(AccImpl->MRange[0]),
+                              &(AccImpl->MAccessRange[0]),
                               sizeof(size_t) * AccImpl->MDims, NextArgId + 1));
           MArgs.emplace_back(
               detail::ArgDesc(detail::kernel_param_kind_t::kind_std_layout,
-                              &AccImpl->MOrigRange[0],
+                              &AccImpl->MMemoryRange[0],
                               sizeof(size_t) * AccImpl->MDims, NextArgId + 2));
           MArgs.emplace_back(
               detail::ArgDesc(detail::kernel_param_kind_t::kind_std_layout,

sycl/include/CL/sycl/id.hpp

@@ -18,6 +18,8 @@ template <int dimensions> class range;
 template <int dimensions = 1> struct id : public detail::array<dimensions> {
 private:
   using base = detail::array<dimensions>;
+  static_assert(dimensions >= 1 && dimensions <= 3,
+                "id can only be 1, 2, or 3 dimentional.");
 public:
   id() = default;
 

sycl/include/CL/sycl/range.hpp

@@ -16,6 +16,8 @@ namespace sycl {
 template <int dimensions> struct id;
 template <int dimensions = 1>
 class range : public detail::array<dimensions> {
+  static_assert(dimensions >= 1 && dimensions <= 3,
+                "range can only be 1, 2, or 3 dimentional.");
   using base = detail::array<dimensions>;
 public:
   /* The following constructor is only available in the range class