[SYCL][USM] Initial commit of USM fill operation (#2305)

Implements the fill operation for USM as defined in SYCL 2020 provisional.

Signed-off-by: James Brodman <[email protected]>

Author: jbrodman

sycl/include/CL/sycl/handler.hpp

@@ -38,6 +38,8 @@ template <typename DataT, int Dimensions, cl::sycl::access::mode AccessMode,
           cl::sycl::access::placeholder IsPlaceholder>
 class __fill;
 
+template <typename T> class __usmfill;
+
 template <typename T_Src, typename T_Dst, int Dims,
           cl::sycl::access::mode AccessMode,
           cl::sycl::access::target AccessTarget,
@@ -1713,6 +1715,22 @@ class __SYCL_EXPORT handler {
     }
   }
 
+  /// Fills the specified memory with the specified pattern.
+  ///
+  /// \param Ptr is the pointer to the memory to fill
+  /// \param Pattern is the pattern to fill into the memory.  T should be
+  /// trivially copyable.
+  /// \param Count is the number of times to fill Pattern into Ptr.
+  template <typename T> void fill(void *Ptr, const T &Pattern, size_t Count) {
+    throwIfActionIsCreated();
+    static_assert(std::is_trivially_copyable<T>::value,
+                  "Pattern must be trivially copyable");
+    parallel_for<class __usmfill<T>>(range<1>(Count), [=](id<1> Index) {
+      T *CastedPtr = static_cast<T *>(Ptr);
+      CastedPtr[Index] = Pattern;
+    });
+  }
+
   /// Prevents any commands submitted afterward to this queue from executing
   /// until all commands previously submitted to this queue have entered the
   /// complete state.

sycl/include/CL/sycl/queue.hpp

@@ -319,6 +319,16 @@ class __SYCL_EXPORT queue {
   /// property, an invalid_object_error SYCL exception.
   template <typename PropertyT> PropertyT get_property() const;
 
+  /// Fills the specified memory with the specified pattern.
+  ///
+  /// \param Ptr is the pointer to the memory to fill
+  /// \param Pattern is the pattern to fill into the memory.  T should be
+  /// trivially copyable.
+  /// \param Count is the number of times to fill Pattern into Ptr.
+  template <typename T> event fill(void *Ptr, const T &Pattern, size_t Count) {
+    return submit([&](handler &CGH) { CGH.fill<T>(Ptr, Pattern, Count); });
+  }
+
   /// Fills the memory pointed by a USM pointer with the value specified.
   /// No operations is done if \param Count is zero. An exception is thrown
   /// if \param Dest is nullptr. The behavior is undefined if \param Ptr

sycl/test/usm/fill.cpp

@@ -0,0 +1,40 @@
+//==---- fill.cpp - USM fill test ------------------------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// XFAIL: cuda
+// piextUSM*Alloc functions for CUDA are not behaving as described in
+//
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple  %s -o %t1.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t1.out
+// RUN: %CPU_RUN_PLACEHOLDER %t1.out
+// RUN: %GPU_RUN_PLACEHOLDER %t1.out
+
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+constexpr int count = 100;
+constexpr int pattern = 42;
+
+int main() {
+  queue q;
+  if (q.get_device().get_info<info::device::usm_shared_allocations>()) {
+    int *mem = malloc_shared<int>(count, q);
+
+    for (int i = 0; i < count; i++)
+      mem[i] = 0;
+
+    q.fill(mem, pattern, count);
+    q.wait();
+
+    for (int i = 0; i < count; i++) {
+      assert(mem[i] == pattern);
+    }
+  }
+  std::cout << "Passed\n";
+  return 0;
+}