[SYCL] Fix accessor fill for unsupported patterns (#8845)

Currently most of the PI backends have similar pattern limitations as
OpenCL, i.e. only certain sizes of patterns are supported. This is not
currently being considered by the implementation however, potentially
causing the backend calls to fail. This commit makes the runtime fall
back to a fill kernel if the pattern has an unsupported size.

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

Author: steffenlarsen

sycl/include/sycl/handler.hpp

@@ -913,6 +913,12 @@ class __SYCL_EXPORT handler {
            AccessMode == access::mode::discard_read_write;
   }
 
+  // PI APIs only support select fill sizes: 1, 2, 4, 8, 16, 32, 64, 128
+  constexpr static bool isBackendSupportedFillSize(size_t Size) {
+    return Size == 1 || Size == 2 || Size == 4 || Size == 8 || Size == 16 ||
+           Size == 32 || Size == 64 || Size == 128;
+  }
+
   template <int Dims, typename LambdaArgType> struct TransformUserItemType {
     using type = typename std::conditional<
         std::is_convertible<nd_item<Dims>, LambdaArgType>::value, nd_item<Dims>,
@@ -2392,15 +2398,17 @@ class __SYCL_EXPORT handler {
   fill(accessor<T, Dims, AccessMode, AccessTarget, IsPlaceholder, PropertyListT>
            Dst,
        const T &Pattern) {
+    assert(!MIsHost && "fill() should no longer be callable on a host device.");
+
     if (Dst.is_placeholder())
       checkIfPlaceholderIsBoundToHandler(Dst);
 
     throwIfActionIsCreated();
     // TODO add check:T must be an integral scalar value or a SYCL vector type
     static_assert(isValidTargetForExplicitOp(AccessTarget),
                   "Invalid accessor target for the fill method.");
-    if (!MIsHost && (((Dims == 1) && isConstOrGlobal(AccessTarget)) ||
-                     isImageOrImageArray(AccessTarget))) {
+    if constexpr (isBackendSupportedFillSize(sizeof(T)) &&
+                  (Dims == 1 || isImageOrImageArray(AccessTarget))) {
       setType(detail::CG::Fill);
 
       detail::AccessorBaseHost *AccBase = (detail::AccessorBaseHost *)&Dst;
@@ -2414,9 +2422,6 @@ class __SYCL_EXPORT handler {
       auto PatternPtr = reinterpret_cast<T *>(MPattern.data());
       *PatternPtr = Pattern;
     } else {
-
-      // TODO: Temporary implementation for host. Should be handled by memory
-      // manger.
       range<Dims> Range = Dst.get_range();
       parallel_for<
           class __fill<T, Dims, AccessMode, AccessTarget, IsPlaceholder>>(

sycl/test-e2e/Basic/fill_accessor.cpp

@@ -0,0 +1,88 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <algorithm>
+#include <array>
+#include <numeric>
+
+using namespace sycl;
+
+size_t NumErrors = 0;
+
+template <typename T, size_t N>
+std::ostream &operator<<(std::ostream &OS, const std::array<T, N> &Arr) {
+  OS << "{";
+  for (size_t I = 0; I < N; ++I) {
+    if (I)
+      OS << ",";
+    OS << Arr[I];
+  }
+  OS << "}";
+  return OS;
+}
+
+template <typename T, int Dims>
+void CheckFill(queue &Q, range<Dims> Range, T Init, T Expected) {
+  std::vector<T> Data(Range.size(), Init);
+  {
+    buffer<T, Dims> Buffer(Data.data(), Range);
+    Q.submit([&](handler &CGH) {
+       accessor Accessor(Buffer, CGH, write_only);
+       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) {
+  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);
+}
+
+int main() {
+  queue Q;
+
+  // Test different power-of-two sizes.
+  CheckFillDifferentDims<char>(Q, 10, 'a', 'z');
+  CheckFillDifferentDims<std::array<char, 2>>(Q, 10, {'a', 'z'}, {'z', 'a'});
+  CheckFillDifferentDims<short>(Q, 10, 8, -16);
+  CheckFillDifferentDims<float>(Q, 10, 123.4, 3.14);
+  CheckFillDifferentDims<uint64_t>(Q, 10, 42, 24);
+  CheckFillDifferentDims<std::array<uint64_t, 2>>(Q, 10, {4, 42}, {24, 4});
+  CheckFillDifferentDims<std::array<uint64_t, 4>>(Q, 10, {4, 42, 424, 4242},
+                                                  {2424, 424, 24, 4});
+  CheckFillDifferentDims<std::array<uint64_t, 8>>(
+      Q, 10, {4, 42, 424, 4242, 42424, 424242, 4242424, 42424242},
+      {24242424, 2424242, 242424, 24242, 2424, 424, 24, 4});
+  CheckFillDifferentDims<std::array<uint64_t, 16>>(
+      Q, 10,
+      {24242424, 2424242, 242424, 24242, 2424, 424, 24, 4, 4, 42, 424, 4242,
+       42424, 424242, 4242424, 42424242},
+      {4, 42, 424, 4242, 42424, 424242, 4242424, 42424242, 24242424, 2424242,
+       242424, 24242, 2424, 424, 24, 4});
+
+  // Test with non-power-of-two sizes.
+  CheckFillDifferentDims<std::array<char, 5>>(Q, 10, {'a', 'b', 'c', 'd', 'e'},
+                                              {'A', 'B', 'C', 'D', 'E'});
+  std::array<char, 129> InitCharArray129;
+  std::fill(InitCharArray129.begin(), InitCharArray129.end(), 130);
+  std::array<char, 129> ExpectedCharArray129;
+  std::iota(ExpectedCharArray129.begin(), ExpectedCharArray129.end(), 1);
+  CheckFillDifferentDims<std::array<char, 129>>(Q, 10, InitCharArray129,
+                                                ExpectedCharArray129);
+
+  return NumErrors;
+}