From 7de840be8f1343002f483d8e02a67d212eeb49e2 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Thu, 9 Jun 2022 06:33:46 -0700
Subject: [PATCH] [SYCL] Replace std::fill in single_task with a fill command

Some targets may not support std::fill in kernel code. This commit
changes a number of std::fill in kernel code to be a fill command on
queue.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 SYCL/Reduction/reduction_nd_N_queue_shortcut.cpp  | 12 +++++-------
 SYCL/Reduction/reduction_nd_queue_shortcut.cpp    |  4 ++--
 SYCL/Reduction/reduction_range_queue_shortcut.cpp |  4 ++--
 3 files changed, 9 insertions(+), 11 deletions(-)

diff --git a/SYCL/Reduction/reduction_nd_N_queue_shortcut.cpp b/SYCL/Reduction/reduction_nd_N_queue_shortcut.cpp
index d7bfcc3b4c..c05287ec9e 100644
--- a/SYCL/Reduction/reduction_nd_N_queue_shortcut.cpp
+++ b/SYCL/Reduction/reduction_nd_N_queue_shortcut.cpp
@@ -74,11 +74,9 @@ int test(queue &Q, BOpT1 BOp1, BOpT2 BOp2, const nd_range<Dims> &Range) {
                    })
         .wait();
   } else if constexpr (TC == TestCase::Dependency) {
-    auto E = Q.single_task([=]() {
-      std::fill(Arr1, Arr1 + NElems, 1);
-      std::fill(Arr2, Arr2 + NElems, 2);
-    });
-    Q.parallel_for(Range, E, R1, R2,
+    auto E1 = Q.fill<T1>(Arr1, 1, NElems);
+    auto E2 = Q.fill<T2>(Arr2, 2, NElems, E1);
+    Q.parallel_for(Range, E2, R1, R2,
                    [=](nd_item<Dims> It, auto &Sum1, auto &Sum2) {
                      size_t LinId = It.get_global_linear_id();
                      Sum1.combine(static_cast<T1>(LinId) + Arr1[LinId]);
@@ -86,8 +84,8 @@ int test(queue &Q, BOpT1 BOp1, BOpT2 BOp2, const nd_range<Dims> &Range) {
                    })
         .wait();
   } else {
-    auto E1 = Q.single_task([=]() { std::fill(Arr1, Arr1 + NElems, 1); });
-    auto E2 = Q.single_task([=]() { std::fill(Arr2, Arr2 + NElems, 2); });
+    auto E1 = Q.fill<T1>(Arr1, 1, NElems);
+    auto E2 = Q.fill<T2>(Arr2, 2, NElems);
     std::vector<event> EVec{E1, E2};
     Q.parallel_for(Range, EVec, R1, R2,
                    [=](nd_item<Dims> It, auto &Sum1, auto &Sum2) {
diff --git a/SYCL/Reduction/reduction_nd_queue_shortcut.cpp b/SYCL/Reduction/reduction_nd_queue_shortcut.cpp
index 904a3900d9..ab0930922d 100644
--- a/SYCL/Reduction/reduction_nd_queue_shortcut.cpp
+++ b/SYCL/Reduction/reduction_nd_queue_shortcut.cpp
@@ -51,13 +51,13 @@ int test(queue &Q, BinaryOperation BOp, const nd_range<Dims> &Range) {
        Sum.combine(static_cast<T>(LinId) + Arr[LinId]);
      }).wait();
   } else if constexpr (TC == TestCase::Dependency) {
-    auto E = Q.single_task([=]() { std::fill(Arr, Arr + NElems, 1); });
+    auto E = Q.fill<T>(Arr, 1, NElems);
     Q.parallel_for(Range, E, Redu, [=](nd_item<Dims> It, auto &Sum) {
        size_t LinId = It.get_global_linear_id();
        Sum.combine(static_cast<T>(LinId) + Arr[LinId]);
      }).wait();
   } else {
-    auto E = Q.single_task([=]() { std::fill(Arr, Arr + NElems, 1); });
+    auto E = Q.fill<T>(Arr, 1, NElems);
     std::vector<event> EVec{E};
     Q.parallel_for(Range, EVec, Redu, [=](nd_item<Dims> It, auto &Sum) {
        size_t LinId = It.get_global_linear_id();
diff --git a/SYCL/Reduction/reduction_range_queue_shortcut.cpp b/SYCL/Reduction/reduction_range_queue_shortcut.cpp
index aecca9067a..79dc520d7b 100644
--- a/SYCL/Reduction/reduction_range_queue_shortcut.cpp
+++ b/SYCL/Reduction/reduction_range_queue_shortcut.cpp
@@ -59,13 +59,13 @@ int test(queue &Q, BinaryOperation BOp, const range<Dims> &Range) {
        Sum.combine(static_cast<T>(LinId) + Arr[LinId]);
      }).wait();
   } else if constexpr (TC == TestCase::Dependency) {
-    auto E = Q.single_task([=]() { std::fill(Arr, Arr + NElems, 1); });
+    auto E = Q.fill<T>(Arr, 1, NElems);
     Q.parallel_for(Range, E, Redu, [=](id<Dims> Id, auto &Sum) {
        size_t LinId = linearizeId(Id, Range);
        Sum.combine(static_cast<T>(LinId) + Arr[LinId]);
      }).wait();
   } else {
-    auto E = Q.single_task([=]() { std::fill(Arr, Arr + NElems, 1); });
+    auto E = Q.fill<T>(Arr, 1, NElems);
     std::vector<event> EVec{E};
     Q.parallel_for(Range, EVec, Redu, [=](id<Dims> Id, auto &Sum) {
        size_t LinId = linearizeId(Id, Range);