[SYCL] Fix empty zero-dimensional accessor access range (#10156)

The current implementation of zero-dimensional accessors always assume
that its size is 1, but this is not the case if the underlying buffer is
empty. This commit fixes this behavior by correcting the used access
range for the accessor, making it correctly report the right sizes and
iterators.

---------

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/include/sycl/accessor.hpp

@@ -276,6 +276,13 @@ template <> struct IsCxPropertyList<ext::oneapi::accessor_property_list<>> {
   constexpr static bool value = false;
 };
 
+// Zero-dimensional accessors references at-most a single element, so the range
+// is either 0 if the associated buffer is empty or 1 otherwise.
+template <typename BufferT>
+sycl::range<1> GetZeroDimAccessRange(BufferT Buffer) {
+  return std::min(Buffer.size(), size_t{1});
+}
+
 __SYCL_EXPORT device getDeviceFromHandler(handler &CommandGroupHandlerRef);
 
 template <typename DataT, int Dimensions, access::mode AccessMode,
@@ -1512,11 +1519,14 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
       const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
 #ifdef __SYCL_DEVICE_ONLY__
-      : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
+      : impl(id<AdjustedDim>(), detail::GetZeroDimAccessRange(BufferRef),
+             BufferRef.get_range()) {
     (void)PropertyList;
 #else
       : AccessorBaseHost(
-            /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
+            /*Offset=*/{0, 0, 0},
+            detail::convertToArrayOfN<3, 1>(
+                detail::GetZeroDimAccessRange(BufferRef)),
             detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
             getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
@@ -1548,11 +1558,14 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
 #ifdef __SYCL_DEVICE_ONLY__
-      : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
+      : impl(id<AdjustedDim>(), detail::GetZeroDimAccessRange(BufferRef),
+             BufferRef.get_range()) {
     (void)PropertyList;
 #else
       : AccessorBaseHost(
-            /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
+            /*Offset=*/{0, 0, 0},
+            detail::convertToArrayOfN<3, 1>(
+                detail::GetZeroDimAccessRange(BufferRef)),
             detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
             getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
@@ -1579,13 +1592,16 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
       const property_list &PropertyList = {},
       const detail::code_location CodeLoc = detail::code_location::current())
 #ifdef __SYCL_DEVICE_ONLY__
-      : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
+      : impl(id<AdjustedDim>(), detail::GetZeroDimAccessRange(BufferRef),
+             BufferRef.get_range()) {
     (void)CommandGroupHandler;
     (void)PropertyList;
   }
 #else
       : AccessorBaseHost(
-            /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
+            /*Offset=*/{0, 0, 0},
+            detail::convertToArrayOfN<3, 1>(
+                detail::GetZeroDimAccessRange(BufferRef)),
             detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
             getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
@@ -1612,13 +1628,16 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
           {},
       const detail::code_location CodeLoc = detail::code_location::current())
 #ifdef __SYCL_DEVICE_ONLY__
-      : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
+      : impl(id<AdjustedDim>(), detail::GetZeroDimAccessRange(BufferRef),
+             BufferRef.get_range()) {
     (void)CommandGroupHandler;
     (void)PropertyList;
   }
 #else
       : AccessorBaseHost(
-            /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
+            /*Offset=*/{0, 0, 0},
+            detail::convertToArrayOfN<3, 1>(
+                detail::GetZeroDimAccessRange(BufferRef)),
             detail::convertToArrayOfN<3, 1>(BufferRef.get_range()),
             getAdjustedMode(PropertyList),
             detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
@@ -2416,10 +2435,7 @@ 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());
+    return detail::convertToArrayOfN<AdjustedDim, 1>(getAccessRange());
   }
 
   template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0)>>

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

@@ -0,0 +1,74 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Disabled for HIP (https://github.com/intel/llvm/issues/10358) and CUDA
+// (https://github.com/intel/llvm/issues/10360)
+// UNSUPPORTED: cuda || hip
+
+// Tests the size and iterator members of an empty zero-dimensional accessor.
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+int check_host(bool CheckResult, std::string Msg) {
+  if (!CheckResult)
+    std::cout << "Case failed: " << Msg << std::endl;
+  return !CheckResult;
+}
+
+int main() {
+  int Failures = 0;
+  queue Q;
+
+  buffer<int, 1> EmptyBuf{0};
+  assert(EmptyBuf.size() == 0);
+
+  {
+    host_accessor<int, 0> EmptyHostAcc{EmptyBuf};
+    Failures += check_host(EmptyHostAcc.empty(), "empty() on host_accesor");
+    Failures += check_host(EmptyHostAcc.size() == 0, "size() on host_accesor");
+    Failures += check_host(EmptyHostAcc.byte_size() == 0,
+                           "byte_size() on host_accesor");
+    Failures +=
+        check_host(EmptyHostAcc.max_size() == 0, "max_size() on host_accesor");
+    Failures += check_host(EmptyHostAcc.begin() == EmptyHostAcc.end(),
+                           "begin()/end() on host_accesor");
+    Failures += check_host(EmptyHostAcc.cbegin() == EmptyHostAcc.cend(),
+                           "cbegin()/cend() on host_accesor");
+    Failures += check_host(EmptyHostAcc.rbegin() == EmptyHostAcc.rend(),
+                           "rbegin()/rend() on host_accesor");
+    Failures += check_host(EmptyHostAcc.crbegin() == EmptyHostAcc.crend(),
+                           "crbegin()/crend() on host_accesor");
+  }
+
+  bool DeviceResults[8] = {false};
+  {
+    buffer<bool, 1> DeviceResultsBuf{DeviceResults, range<1>{8}};
+    Q.submit([&](handler &CGH) {
+      accessor<int, 0> EmptyDevAcc{EmptyBuf, CGH};
+      accessor DeviceResultsAcc{DeviceResultsBuf, CGH};
+      CGH.single_task([=]() {
+        DeviceResultsAcc[0] = EmptyDevAcc.empty();
+        DeviceResultsAcc[1] = EmptyDevAcc.size() == 0;
+        DeviceResultsAcc[2] = EmptyDevAcc.byte_size() == 0;
+        DeviceResultsAcc[3] = EmptyDevAcc.max_size() == 0;
+        DeviceResultsAcc[4] = EmptyDevAcc.begin() == EmptyDevAcc.end();
+        DeviceResultsAcc[5] = EmptyDevAcc.cbegin() == EmptyDevAcc.cend();
+        DeviceResultsAcc[6] = EmptyDevAcc.rbegin() == EmptyDevAcc.rend();
+        DeviceResultsAcc[7] = EmptyDevAcc.crbegin() == EmptyDevAcc.crend();
+      });
+    });
+  }
+
+  Failures += check_host(DeviceResults[0], "empty() on accessor");
+  Failures += check_host(DeviceResults[1], "size() on accessor");
+  Failures += check_host(DeviceResults[2], "byte_size() on accessor");
+  Failures += check_host(DeviceResults[3], "max_size() on accessor");
+  Failures += check_host(DeviceResults[4], "begin()/end() on accessor");
+  Failures += check_host(DeviceResults[5], "cbegin()/cend() on accessor");
+  Failures += check_host(DeviceResults[6], "rbegin()/rend() on accessor");
+  Failures += check_host(DeviceResults[7], "crbegin()/crend() on accessor");
+
+  return Failures;
+}