diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 0985e60d8eeb9..55b08564549a6 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -905,6 +905,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>,
@@ -2384,6 +2390,8 @@ 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);
 
@@ -2391,8 +2399,8 @@ class __SYCL_EXPORT handler {
     // 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;
@@ -2406,9 +2414,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>>(
diff --git a/sycl/test-e2e/Basic/fill_accessor.cpp b/sycl/test-e2e/Basic/fill_accessor.cpp
new file mode 100644
index 0000000000000..96219da67ea28
--- /dev/null
+++ b/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;
+}