[SYCL] Fix bug in fill op for zero-dim accessor. (#9287)

According to the SYCL2020 spec, a zero-dim accessor is used for
accessing the first element of a buffer. Also, `range` and `id` classes
cannot accept zero dimensions. Thus, we need a special case for zero-dim
accessors that simply fills the first element.

Closes #9252.

---------

Signed-off-by: Maronas, Marcos <[email protected]>

Author: maarquitos14

sycl/include/sycl/handler.hpp

@@ -2432,7 +2432,7 @@ class __SYCL_EXPORT handler {
     static_assert(isValidTargetForExplicitOp(AccessTarget),
                   "Invalid accessor target for the fill method.");
     if constexpr (isBackendSupportedFillSize(sizeof(T)) &&
-                  (Dims == 1 || isImageOrImageArray(AccessTarget))) {
+                  (Dims <= 1 || isImageOrImageArray(AccessTarget))) {
       setType(detail::CG::Fill);
 
       detail::AccessorBaseHost *AccBase = (detail::AccessorBaseHost *)&Dst;
@@ -2445,6 +2445,11 @@ class __SYCL_EXPORT handler {
       MPattern.resize(sizeof(T));
       auto PatternPtr = reinterpret_cast<T *>(MPattern.data());
       *PatternPtr = Pattern;
+    } else if constexpr (Dims == 0) {
+      // Special case for zero-dim accessors.
+      parallel_for<
+          class __fill<T, Dims, AccessMode, AccessTarget, IsPlaceholder>>(
+          range<1>(1), [=](id<1> Index) { Dst = Pattern; });
     } else {
       range<Dims> Range = Dst.get_range();
       parallel_for<

sycl/source/detail/memory_manager.cpp

@@ -781,7 +781,7 @@ void MemoryManager::fill(SYCLMemObjI *SYCLMemObj, void *Mem, QueueImplPtr Queue,
 
   const detail::plugin &Plugin = Queue->getPlugin();
   if (SYCLMemObj->getType() == detail::SYCLMemObjI::MemObjType::Buffer) {
-    if (Dim == 1) {
+    if (Dim <= 1) {
       Plugin.call<PiApiKind::piEnqueueMemBufferFill>(
           Queue->getHandleRef(), pi::cast<RT::PiMem>(Mem), Pattern, PatternSize,
           Offset[0] * ElementSize, Range[0] * ElementSize, DepEvents.size(),

sycl/test-e2e/Basic/fill_accessor.cpp

@@ -45,8 +45,31 @@ void CheckFill(queue &Q, range<Dims> Range, T Init, T Expected) {
   }
 }
 
+template <typename T>
+void CheckFillZeroDimAccessor(queue &Q, T Init, T Expected) {
+  constexpr int Dims = 1;
+  range<1> Range(1);
+  std::vector<T> Data(Range.size(), Init);
+  {
+    buffer<T, Dims> Buffer(Data.data(), Range);
+    Q.submit([&](handler &CGH) {
+       accessor<T, 0, sycl::access::mode::write> Accessor(Buffer, CGH);
+       CGH.fill(Accessor, Expected);
+     }).wait_and_throw();
+  }
+  for (size_t I = 0; I < Range.size(); ++I) {
+    if (Data[I] != Expected) {
+      std::cout << "Unexpected value " << Data[I] << " at index " << I
+                << " after fill. Expected " << Expected << "." << std::endl;
+      ++NumErrors;
+      return;
+    }
+  }
+}
+
 template <typename T>
 void CheckFillDifferentDims(queue &Q, size_t N, T Init, T Expected) {
+  CheckFillZeroDimAccessor<T>(Q, Init, Expected);
   CheckFill<T>(Q, range<1>{N}, Init, Expected);
   CheckFill<T>(Q, range<2>{N, N}, Init, Expected);
   CheckFill<T>(Q, range<3>{N, N, N}, Init, Expected);