[SYCL] Support of iterator operations for zero-dimension accessors. (#9141)

* Zero dimension accessors allow access to only one element of the
corresponding buffer. This PR handles iterator operators for zero
dimension accessors by limiting the supported range to 1 element using
`AccessorRange`.
* Add support for 0-dim `local_accessor` iterator functionality. 
* Adds relevant tests.

Author: mmoadeli

sycl/include/sycl/accessor.hpp

@@ -1255,9 +1255,9 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
       std::conditional_t<AccessTarget == access::target::device,
                          global_ptr<value_type, IsDecorated>, value_type *>;
 
-  using iterator = typename detail::accessor_iterator<value_type, Dimensions>;
+  using iterator = typename detail::accessor_iterator<value_type, AdjustedDim>;
   using const_iterator =
-      typename detail::accessor_iterator<const value_type, Dimensions>;
+      typename detail::accessor_iterator<const value_type, AdjustedDim>;
   using reverse_iterator = std::reverse_iterator<iterator>;
   using const_reverse_iterator = std::reverse_iterator<const_iterator>;
   using difference_type =
@@ -2035,18 +2035,16 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
 
   bool empty() const noexcept { return size() == 0; }
 
-  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
+  template <int Dims = Dimensions,
+            typename = std::enable_if_t<Dims == Dimensions && (Dims > 0)>>
   range<Dimensions> get_range() const {
-    return detail::convertToArrayOfN<Dimensions, 1>(getAccessRange());
+    return getRange<Dims>();
   }
 
-  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
+  template <int Dims = Dimensions,
+            typename = std::enable_if_t<Dims == Dimensions && (Dims > 0)>>
   id<Dimensions> get_offset() const {
-    static_assert(
-        !(PropertyListT::template has_property<
-            sycl::ext::oneapi::property::no_offset>()),
-        "Accessor has no_offset property, get_offset() can not be used");
-    return detail::convertToArrayOfN<Dimensions, 0>(getOffset());
+    return getOffset<Dims>();
   }
 
   template <int Dims = Dimensions, typename RefT = RefType,
@@ -2197,29 +2195,29 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
   iterator begin() const noexcept {
     return iterator::getBegin(
         get_pointer(),
-        detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
-        get_offset());
+        detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
+        getRange<AdjustedDim>(), getOffset<AdjustedDim>());
   }
 
   iterator end() const noexcept {
     return iterator::getEnd(
         get_pointer(),
-        detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
-        get_offset());
+        detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
+        getRange<AdjustedDim>(), getOffset<AdjustedDim>());
   }
 
   const_iterator cbegin() const noexcept {
     return const_iterator::getBegin(
         get_pointer(),
-        detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
-        get_offset());
+        detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
+        getRange<AdjustedDim>(), getOffset<AdjustedDim>());
   }
 
   const_iterator cend() const noexcept {
     return const_iterator::getEnd(
         get_pointer(),
-        detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
-        get_offset());
+        detail::convertToArrayOfN<AdjustedDim, 1>(getMemoryRange()),
+        getRange<AdjustedDim>(), getOffset<AdjustedDim>());
   }
 
   reverse_iterator rbegin() const noexcept { return reverse_iterator(end()); }
@@ -2233,6 +2231,23 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
   }
 
 private:
+  template <int Dims, typename = std::enable_if_t<(Dims > 0)>>
+  range<Dims> getRange() const {
+    if constexpr (Dimensions == 0)
+      return range<1>{1};
+    else
+      return detail::convertToArrayOfN<Dims, 1>(getAccessRange());
+  }
+
+  template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>
+  id<Dims> getOffset() const {
+    static_assert(
+        !(PropertyListT::template has_property<
+            sycl::ext::oneapi::property::no_offset>()),
+        "Accessor has no_offset property, get_offset() can not be used");
+    return detail::convertToArrayOfN<Dims, 0>(getOffset());
+  }
+
 #ifdef __SYCL_DEVICE_ONLY__
   size_t getTotalOffset() const noexcept {
     size_t TotalOffset = 0;
@@ -2844,9 +2859,17 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
   bool empty() const noexcept { return this->size() == 0; }
 
   iterator begin() const noexcept {
-    return &this->operator[](id<Dimensions>());
+    if constexpr (Dimensions == 0)
+      return local_acc::getQualifiedPtr();
+    else
+      return &this->operator[](id<Dimensions>());
+  }
+  iterator end() const noexcept {
+    if constexpr (Dimensions == 0)
+      return begin() + 1;
+    else
+      return begin() + this->size();
   }
-  iterator end() const noexcept { return begin() + this->size(); }
 
   const_iterator cbegin() const noexcept { return const_iterator(begin()); }
   const_iterator cend() const noexcept { return const_iterator(end()); }

sycl/test-e2e/Basic/accessor/accessor.cpp

@@ -999,6 +999,21 @@ int main() {
     }
     assert(vec1[7] == 4 && vec2[15] == 4);
   }
+
+  // 0-dim host_accessor iterator
+  {
+    std::vector<int> vec1(8);
+    {
+      sycl::buffer<int> buf1(vec1.data(), vec1.size());
+      sycl::host_accessor<int, 0> acc1(buf1);
+      *acc1.begin() = 4;
+      auto value = *acc1.cbegin();
+      value += *acc1.crbegin();
+      *acc1.rbegin() += value;
+    }
+    assert(vec1[0] == 12);
+  }
+
   // Test swap() on basic accessor
   {
     std::vector<int> vec1(8), vec2(16);
@@ -1095,6 +1110,59 @@ int main() {
     assert(Data == 64);
   }
 
+  // iterator operations test for 0-dim buffer accessor
+  {
+    sycl::queue Queue;
+    int Data[] = {32, 32};
+
+    // Explicit block to prompt copy-back to Data
+    {
+      sycl::buffer<int, 1> DataBuffer(Data, sycl::range<1>(2));
+
+      Queue.submit([&](sycl::handler &CGH) {
+        sycl::accessor<int, 0> Acc(DataBuffer, CGH);
+        CGH.single_task<class acc_0_dim_iter_assignment>([=]() {
+          *Acc.begin() = 64;
+          auto value = *Acc.cbegin();
+          value += *Acc.crbegin();
+          *Acc.rbegin() += value;
+        });
+      });
+      Queue.wait();
+    }
+
+    assert(Data[0] == 64 * 3);
+    assert(Data[1] == 32);
+  }
+
+  // iterator operations test for 0-dim buffer accessor with target::host_task
+  {
+    sycl::queue Queue;
+    int Data[] = {32, 32};
+
+    using HostTaskAcc = sycl::accessor<int, 0, sycl::access::mode::read_write,
+                                       sycl::access::target::host_task>;
+
+    // Explicit block to prompt copy-back to Data
+    {
+      sycl::buffer<int, 1> DataBuffer(Data, sycl::range<1>(2));
+
+      Queue.submit([&](sycl::handler &CGH) {
+        HostTaskAcc Acc(DataBuffer, CGH);
+        CGH.host_task([=]() {
+          *Acc.begin() = 64;
+          auto value = *Acc.cbegin();
+          value += *Acc.crbegin();
+          *Acc.rbegin() += value;
+        });
+      });
+      Queue.wait();
+    }
+
+    assert(Data[0] == 64 * 3);
+    assert(Data[1] == 32);
+  }
+
   // Assignment operator test for 0-dim local accessor
   {
     sycl::queue Queue;
@@ -1171,6 +1239,32 @@ int main() {
     }
   }
 
+  // Assignment operator test for 0-dim local accessor iterator
+  {
+    sycl::queue Queue;
+    int Data = 0;
+
+    // Explicit block to prompt copy-back to Data
+    {
+      sycl::buffer<int, 1> DataBuffer(&Data, sycl::range<1>(1));
+
+      Queue.submit([&](sycl::handler &CGH) {
+        sycl::accessor<int, 0> Acc(DataBuffer, CGH);
+        sycl::local_accessor<int, 0> LocalAcc(CGH);
+        CGH.parallel_for<class local_acc_0_dim_iter_assignment>(
+            sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1> ID) {
+              *LocalAcc.begin() = 32;
+              auto value = *LocalAcc.cbegin();
+              value += *LocalAcc.crbegin();
+              *LocalAcc.rbegin() += value;
+              Acc = LocalAcc;
+            });
+      });
+    }
+
+    assert(Data == 96);
+  }
+
   // host_accessor hash
   {
     sycl::buffer<int> buffer1{sycl::range<1>{1}};