[SYCL] Add LIT test for (reduction + range + USM) cases (#398)

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

Author: v-klochkov

SYCL/Reduction/reduction_range_scalar.hpp

@@ -7,17 +7,28 @@
 using namespace cl::sycl;
 
 template <typename T, bool B> class KName;
+template <typename T, typename> class TName;
 
 template <typename Name, bool IsSYCL2020, access::mode Mode, int AccDim = 1,
           typename T, class BinaryOperation, int Dims>
 int test(queue &Q, T Identity, T Init, BinaryOperation BOp,
          const range<Dims> &Range) {
   printTestLabel<T, BinaryOperation>(IsSYCL2020, Range);
 
-  // Skip the test for such big arrays now.
-  constexpr size_t TwoGB = 2LL * 1024 * 1024 * 1024;
-  if (Range.size() > TwoGB)
+  // It is a known problem with passing data that is close to 4Gb in size
+  // to device. Such data breaks the execution pretty badly.
+  // Some of test cases calling this function try to verify the correctness
+  // of reduction with the global range bigger than the maximal work-group size
+  // for the device. Maximal WG size for device may be very big, e.g. it is
+  // 67108864 for ACC emulator. Multiplying that by some factor
+  // (to exceed max WG-Size) and multiplying it by the element size may exceed
+  // the safe size of data passed to device.
+  // Let's set it to 1 GB for now, and just skip the test if it exceeds 1Gb.
+  constexpr size_t OneGB = 1LL * 1024 * 1024 * 1024;
+  if (Range.size() * sizeof(T) > OneGB) {
+    std::cout << " SKIPPED due to too big data size" << std::endl;
     return 0;
+  }
 
   buffer<T, Dims> InBuf(Range);
   buffer<T, 1> OutBuf(1);
@@ -55,3 +66,115 @@ int testBoth(queue &Q, T Identity, T Init, BinaryOperation BOp,
   return test<KName<Name, false>, false, Mode>(Q, Identity, Init, BOp, Range) +
          test<KName<Name, true>, true, Mode>(Q, Identity, Init, BOp, Range);
 }
+
+template <typename Name, bool IsSYCL2020, access::mode Mode, typename T,
+          class BinaryOperation, int Dims>
+int testUSM(queue &Q, T Identity, T Init, BinaryOperation BOp,
+            const range<Dims> &Range, usm::alloc AllocType) {
+  printTestLabel<T, BinaryOperation>(IsSYCL2020, Range);
+
+  auto Dev = Q.get_device();
+  if (!Dev.has(getUSMAspect(AllocType))) {
+    std::cout << " SKIPPED due to unsupported USM alloc type" << std::endl;
+    return 0;
+  }
+
+  // It is a known problem with passing data that is close to 4Gb in size
+  // to device. Such data breaks the execution pretty badly.
+  // Some of test cases calling this function try to verify the correctness
+  // of reduction with the global range bigger than the maximal work-group size
+  // for the device. Maximal WG size for device may be very big, e.g. it is
+  // 67108864 for ACC emulator. Multiplying that by some factor
+  // (to exceed max WG-Size) and multiplying it by the element size may exceed
+  // the safe size of data passed to device.
+  // Let's set it to 1 GB for now, and just skip the test if it exceeds 1Gb.
+  constexpr size_t OneGB = 1LL * 1024 * 1024 * 1024;
+  if (Range.size() * sizeof(T) > OneGB) {
+    std::cout << " SKIPPED due to too big data size" << std::endl;
+    return 0;
+  }
+
+  T *ReduVarPtr = (T *)malloc(sizeof(T), Dev, Q.get_context(), AllocType);
+  if (ReduVarPtr == nullptr) {
+    std::cout << " SKIPPED due to unrelated reason: alloc returned nullptr"
+              << std::endl;
+    return 0;
+  }
+  if (AllocType == usm::alloc::device) {
+    Q.submit([&](handler &CGH) {
+       CGH.single_task<TName<Name, class InitKernel>>(
+           [=]() { *ReduVarPtr = Init; });
+     }).wait();
+  } else {
+    *ReduVarPtr = Init;
+  }
+
+  // Initialize.
+  T CorrectOut;
+  buffer<T, Dims> InBuf(Range);
+  initInputData(InBuf, CorrectOut, Identity, BOp, Range);
+  if constexpr (Mode == access::mode::read_write)
+    CorrectOut = BOp(CorrectOut, Init);
+
+  // Compute.
+  Q.submit([&](handler &CGH) {
+     auto In = InBuf.template get_access<access::mode::read>(CGH);
+     auto Redu = createReduction<IsSYCL2020, Mode>(ReduVarPtr, Identity, BOp);
+     CGH.parallel_for<TName<Name, class Test>>(
+         Range, Redu, [=](id<Dims> Id, auto &Sum) { Sum.combine(In[Id]); });
+   }).wait();
+
+  // Check correctness.
+  T ComputedOut;
+  if (AllocType == usm::alloc::device) {
+    buffer<T, 1> Buf(&ComputedOut, range<1>(1));
+    Q.submit([&](handler &CGH) {
+       auto OutAcc = Buf.template get_access<access::mode::discard_write>(CGH);
+       CGH.single_task<TName<Name, class Check>>(
+           [=]() { OutAcc[0] = *ReduVarPtr; });
+     }).wait();
+    ComputedOut = (Buf.template get_access<access::mode::read>())[0];
+  } else {
+    ComputedOut = *ReduVarPtr;
+  }
+
+  std::string AllocStr =
+      "AllocMode=" + std::to_string(static_cast<int>(AllocType));
+  int Error = checkResults(Q, IsSYCL2020, BOp, Range, ComputedOut, CorrectOut,
+                           AllocStr);
+  free(ReduVarPtr, Q.get_context());
+  return Error;
+}
+
+template <typename Name, access::mode Mode, typename T, class BinaryOperation,
+          int Dims>
+int test2020USM(queue &Q, T Identity, T Init, BinaryOperation BOp,
+                const range<Dims> &Range) {
+  int NumErrors = 0;
+  NumErrors += testUSM<TName<Name, class Shared2020>, true, Mode, T>(
+      Q, Identity, Init, BOp, Range, usm::alloc::shared);
+  NumErrors += testUSM<TName<Name, class Host2020>, true, Mode, T>(
+      Q, Identity, Init, BOp, Range, usm::alloc::host);
+  NumErrors += testUSM<TName<Name, class Device2020>, true, Mode, T>(
+      Q, Identity, Init, BOp, Range, usm::alloc::device);
+  return NumErrors;
+}
+
+template <typename Name, access::mode Mode, typename T, class BinaryOperation,
+          int Dims>
+int testONEAPIUSM(queue &Q, T Identity, T Init, BinaryOperation BOp,
+                  const range<Dims> &Range) {
+  int NumErrors = 0;
+  if (Mode == access::mode::discard_write) {
+    std::cerr << "Skipped an incorrect test case: ext::oneapi::reduction "
+              << "does not support discard_write mode for USM variables.";
+    return 0;
+  }
+  NumErrors += testUSM<TName<Name, class Shared>, false, Mode, T>(
+      Q, Identity, Init, BOp, Range, usm::alloc::shared);
+  NumErrors += testUSM<TName<Name, class Host>, false, Mode, T>(
+      Q, Identity, Init, BOp, Range, usm::alloc::host);
+  NumErrors += testUSM<TName<Name, class Device>, false, Mode, T>(
+      Q, Identity, Init, BOp, Range, usm::alloc::device);
+  return NumErrors;
+}

SYCL/Reduction/reduction_range_usm_dw.cpp

@@ -0,0 +1,50 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+// TODO: accelerator may not suport atomics required by the current
+// implementation. Enable testing when implementation is fixed.
+// RUNx: %ACC_RUN_PLACEHOLDER %t.out
+
+#include "reduction_range_scalar.hpp"
+
+// This test performs basic checks of parallel_for(range<1>, reduction, func)
+// with reductions initialized with 1-dimensional discard_write accessor
+// accessing 1 element buffer.
+
+using namespace cl::sycl;
+
+int NumErrors = 0;
+
+template <typename Name, typename T, class BinaryOperation, int Dims>
+void tests(queue &Q, T Identity, T Init, BinaryOperation BOp,
+           const range<Dims> &Range) {
+  constexpr access::mode DW = access::mode::discard_write;
+  NumErrors += test2020USM<Name, DW>(Q, Identity, Init, BOp, Range);
+}
+
+int main() {
+  queue Q;
+  printDeviceInfo(Q);
+  size_t MaxWGSize =
+      Q.get_device().get_info<info::device::max_work_group_size>();
+
+  // Fast-reduce and Fast-atomics. Try various range types/sizes.
+  tests<class A1, int>(Q, 0, 99, std::plus<>{}, range<1>{MaxWGSize * 2 + 5});
+  tests<class A2, float>(Q, 0, 99, std::plus<>{}, range<2>{1, 1});
+  tests<class A3, int>(Q, 0, 99, std::plus<>{}, range<3>{MaxWGSize, 1, 2});
+
+  // Try various operations.
+  tests<class B1, int>(Q, ~0, 99, std::bit_and<>{}, range<1>{8});
+  tests<class B2, int>(Q, 0, 0xff99, std::bit_xor<>{}, range<1>{MaxWGSize + 1});
+  tests<class B4, short>(Q, 1, 2, std::multiplies<>{}, range<1>{7});
+  tests<class B5, int>(Q, (std::numeric_limits<int>::max)(), -99,
+                       ext::oneapi::minimum<>{}, range<2>{MaxWGSize, 2});
+
+  // Check with CUSTOM type.
+  tests<class C1>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
+                  CustomVecPlus<long long>{}, range<2>{3, MaxWGSize});
+
+  printFinalStatus(NumErrors);
+  return NumErrors;
+}

SYCL/Reduction/reduction_utils.hpp

@@ -204,13 +204,15 @@ int checkResults(queue &Q, bool IsSYCL2020, BinaryOperation,
     Passed = ComputedRes == CorrectRes;
   }
 
+  if (!AddInfo.empty())
+    AddInfo = std::string(", ") + AddInfo;
   std::cout << AddInfo << (Passed ? ". PASSED" : ". FAILED") << std::endl;
   if (!Passed) {
     printDeviceInfo(Q, true);
     printTestLabel<T, BinaryOperation>(IsSYCL2020, Range, true);
     std::cerr << ", Computed value=" << ComputedRes
-              << ", Expected value=" << CorrectRes << ErrorStr
-              << (AddInfo.empty() ? "" : ", " + AddInfo) << std::endl;
+              << ", Expected value=" << CorrectRes << ErrorStr << AddInfo
+              << std::endl;
   }
   return Passed ? 0 : 1;
 }