[SYCL] Add test cases for SYCL-2020 reductions used in read_write mode

These changes verify intel/llvm#3315

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

Author: v-klochkov

SYCL/Reduction/reduction_nd_N_vars.cpp

@@ -12,6 +12,11 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DTEST_SYCL2020_REDUCTIONS %s -o %t2020.out
+// RUN: %CPU_RUN_PLACEHOLDER %t2020.out
+// RUN: %GPU_RUN_PLACEHOLDER %t2020.out
+// RUN: %ACC_RUN_PLACEHOLDER %t2020.out
+
 // This test checks handling of parallel_for() accepting nd_range and
 // two or more reductions.
 
@@ -125,10 +130,22 @@ int runTest(T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
      auto In3 = InBuf3.template get_access<access::mode::read>(CGH);
      auto In4 = InBuf4.template get_access<access::mode::read>(CGH);
 
+#ifdef TEST_SYCL2020_REDUCTIONS
+     auto Redu1 = sycl::reduction(OutBuf1, CGH, IdentityVal1, BOp1);
+     auto Redu2 = sycl::reduction(OutBuf2, CGH, IdentityVal2, BOp2);
+     auto Redu3 = sycl::reduction(OutBuf3, CGH, IdentityVal3, BOp3);
+     auto Redu4 = sycl::reduction(Out4, IdentityVal4, BOp4);
+#else
      auto Out1 = OutBuf1.template get_access<Mode1>(CGH);
      auto Out2 = OutBuf2.template get_access<Mode2>(CGH);
      accessor<T3, 0, Mode3, access::target::global_buffer> Out3(OutBuf3, CGH);
 
+     auto Redu1 = ONEAPI::reduction(Out1, IdentityVal1, BOp1);
+     auto Redu2 = ONEAPI::reduction(Out2, IdentityVal2, BOp2);
+     auto Redu3 = ONEAPI::reduction(Out3, IdentityVal3, BOp3);
+     auto Redu4 = ONEAPI::reduction(Out4, IdentityVal4, BOp4);
+#endif
+
      auto Lambda = [=](nd_item<1> NDIt, auto &Sum1, auto &Sum2, auto &Sum3,
                        auto &Sum4) {
        size_t I = NDIt.get_global_id(0);
@@ -138,15 +155,6 @@ int runTest(T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
        Sum4.combine(In4[I]);
      };
 
-     auto Redu1 =
-         ONEAPI::reduction<T1, BinaryOperation1>(Out1, IdentityVal1, BOp1);
-     auto Redu2 =
-         ONEAPI::reduction<T2, BinaryOperation2>(Out2, IdentityVal2, BOp2);
-     auto Redu3 =
-         ONEAPI::reduction<T3, BinaryOperation3>(Out3, IdentityVal3, BOp3);
-     auto Redu4 =
-         ONEAPI::reduction<T4, BinaryOperation4>(Out4, IdentityVal4, BOp4);
-
      auto NDR = nd_range<1>{range<1>(NWorkItems), range<1>{WGSize}};
      CGH.parallel_for<ReductionExample>(NDR, Redu1, Redu2, Redu3, Redu4,
                                         Lambda);
@@ -186,17 +194,23 @@ int runTest(T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
 }
 
 int main() {
-  int Error =
-      runTest<class ReduFloatPlus16x1, float, access::mode::discard_write, int,
-              access::mode::read_write, short, access::mode::read_write, int>(
-          0, 1000, std::plus<float>{}, 0, 2000, std::plus<>{}, 0, 4000,
-          std::bit_or<>{}, 0, 8000, std::bit_xor<>{}, usm::alloc::shared, 16,
-          16);
+  constexpr access::mode ReadWriteMode = access::mode::read_write;
+#ifdef TEST_SYCL2020_REDUCTIONS
+  // TODO: property::reduction::initialize_to_identity is not supported yet.
+  // Thus only read_write mode is tested now.
+  constexpr access::mode DiscardWriteMode = access::mode::read_write;
+#else
+  constexpr access::mode DiscardWriteMode = access::mode::discard_write;
+#endif
+
+  int Error = runTest<class ReduFloatPlus16x1, float, DiscardWriteMode, int,
+                      ReadWriteMode, short, ReadWriteMode, int>(
+      0, 1000, std::plus<float>{}, 0, 2000, std::plus<>{}, 0, 4000,
+      std::bit_or<>{}, 0, 8000, std::bit_xor<>{}, usm::alloc::shared, 16, 16);
 
   auto Add = [](auto x, auto y) { return (x + y); };
-  Error += runTest<class ReduFloatPlus5x257, float, access::mode::read_write,
-                   int, access::mode::read_write, short,
-                   access::mode::discard_write, int>(
+  Error += runTest<class ReduFloatPlus5x257, float, ReadWriteMode, int,
+                   ReadWriteMode, short, DiscardWriteMode, int>(
       0, 1000, std::plus<float>{}, 0, 2000, std::plus<>{}, 0, 4000, Add, 0,
       8000, std::bit_xor<int>{}, usm::alloc::device, 5 * (256 + 1), 5);
 

SYCL/Reduction/reduction_nd_ext_double.cpp

@@ -3,8 +3,14 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DTEST_SYCL2020_REDUCTIONS %s -o %t2020.out
+// RUN: %CPU_RUN_PLACEHOLDER %t2020.out
+// RUN: %GPU_RUN_PLACEHOLDER %t2020.out
+// RUN: %ACC_RUN_PLACEHOLDER %t2020.out
+
 // TODO: Enable the test for HOST when it supports intel::reduce() and barrier()
 // RUNx: %HOST_RUN_PLACEHOLDER %t.out
+// RUNx: %HOST_RUN_PLACEHOLDER %t2020.out
 
 // This test performs basic checks of parallel_for(nd_range, reduction, func)
 // used with 'double' type.

SYCL/Reduction/reduction_nd_ext_half.cpp

@@ -1,12 +1,18 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DTEST_SYCL2020_REDUCTIONS %s -o %t2020.out
+// RUN: %GPU_RUN_PLACEHOLDER %t2020.out
+
 // TODO: Enable the test for CPU/ACC when they support half type.
 // RUNx: %CPU_RUN_PLACEHOLDER %t.out
 // RUNx: %ACC_RUN_PLACEHOLDER %t.out
+// RUNx: %CPU_RUN_PLACEHOLDER %t2020.out
+// RUNx: %ACC_RUN_PLACEHOLDER %t2020.out
 
 // TODO: Enable the test for HOST when it supports intel::reduce() and barrier()
 // RUNx: %HOST_RUN_PLACEHOLDER %t.out
+// RUNx: %HOST_RUN_PLACEHOLDER %t2020.out
 
 // This test performs basic checks of parallel_for(nd_range, reduction, func)
 // used with 'half' type.

SYCL/Reduction/reduction_nd_ext_type.hpp

@@ -26,8 +26,12 @@ void test(T Identity, size_t WGSize, size_t NWItems) {
   queue Q;
   Q.submit([&](handler &CGH) {
     auto In = InBuf.template get_access<access::mode::read>(CGH);
+#ifdef TEST_SYCL2020_REDUCTIONS
+    auto Redu = sycl::reduction(OutBuf, CGH, Identity, BOp);
+#else
     accessor<T, Dim, Mode, access::target::global_buffer> Out(OutBuf, CGH);
     auto Redu = ONEAPI::reduction(Out, Identity, BOp);
+#endif
 
     range<1> GlobalRange(NWItems);
     range<1> LocalRange(WGSize);
@@ -59,25 +63,29 @@ template <typename T> int runTests(const string_class &ExtensionName) {
     return 0;
   }
 
+  constexpr access::mode ReadWriteMode = access::mode::read_write;
+#ifdef TEST_SYCL2020_REDUCTIONS
+  // TODO: property::reduction::initialize_to_identity is not supported yet.
+  // Thus only read_write mode is tested now.
+  constexpr access::mode DiscardWriteMode = access::mode::read_write;
+#else
+  constexpr access::mode DiscardWriteMode = access::mode::discard_write;
+#endif
+
   // Check some less standards WG sizes and corner cases first.
-  test<class KernelName_oTh, T, 1, access::mode::read_write,
-       std::multiplies<T>>(0, 4, 4);
-  test<class KernelName_QUQnMARQT, T, 0, access::mode::discard_write,
-       ONEAPI::plus<T>>(0, 4, 64);
+  test<class A, T, 1, ReadWriteMode, std::multiplies<T>>(0, 4, 4);
+  test<class B, T, 0, DiscardWriteMode, ONEAPI::plus<T>>(0, 4, 64);
 
-  test<class KernelName_xGixNo, T, 0, access::mode::read_write,
-       ONEAPI::minimum<T>>(getMaximumFPValue<T>(), 7, 7);
-  test<class KernelName_qXNFw, T, 1, access::mode::discard_write,
-       ONEAPI::maximum<T>>(getMinimumFPValue<T>(), 7, 7 * 5);
+  test<class C, T, 0, ReadWriteMode, ONEAPI::minimum<T>>(getMaximumFPValue<T>(),
+                                                         7, 7);
+  test<class D, T, 1, access::mode::discard_write, ONEAPI::maximum<T>>(
+      getMinimumFPValue<T>(), 7, 7 * 5);
 
-#if __cplusplus >= 201402L
-  test<class KernelName_lXdWtzANdDcvm, T, 1, access::mode::read_write,
-       ONEAPI::plus<>>(1, 3, 3 * 5);
-  test<class KernelName_FDQalsDxmbi, T, 1, access::mode::discard_write,
-       ONEAPI::minimum<>>(getMaximumFPValue<T>(), 3, 3);
-  test<class KernelName_TaNRRxDRXbzYrFImPYC, T, 0, access::mode::discard_write,
-       ONEAPI::maximum<>>(getMinimumFPValue<T>(), 3, 3);
-#endif // __cplusplus >= 201402L
+  test<class E, T, 1, ReadWriteMode, ONEAPI::plus<>>(1, 3, 3 * 5);
+  test<class F, T, 1, DiscardWriteMode, ONEAPI::minimum<>>(
+      getMaximumFPValue<T>(), 3, 3);
+  test<class G, T, 0, DiscardWriteMode, ONEAPI::maximum<>>(
+      getMinimumFPValue<T>(), 3, 3);
 
   std::cout << "Test passed\n";
   return 0;

SYCL/Reduction/reduction_nd_s0_rw.cpp

@@ -1,9 +1,13 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
-// RUNx: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DTEST_SYCL2020_REDUCTIONS %s -o %t2020.out
+// RUN: %CPU_RUN_PLACEHOLDER %t2020.out
+// RUN: %GPU_RUN_PLACEHOLDER %t2020.out
+// RUN: %ACC_RUN_PLACEHOLDER %t2020.out
+
 // This test performs basic checks of parallel_for(nd_range, reduction, func)
 // with reductions initialized with 0-dimensional read_write accessor.
 
@@ -13,11 +17,15 @@
 
 using namespace cl::sycl;
 
+// This allocator is needed only for the purpose of testing buffers
+// with allocator that is not same_as sycl::buffer_allocator.
+struct CustomAllocator : public sycl::buffer_allocator {};
+
 template <typename SpecializationKernelName, typename T, int Dim,
           class BinaryOperation>
 void test(T Identity, size_t WGSize, size_t NWItems) {
   buffer<T, 1> InBuf(NWItems);
-  buffer<T, 1> OutBuf(1);
+  buffer<T, 1, CustomAllocator> OutBuf(1);
 
   // Initialize.
   BinaryOperation BOp;
@@ -30,9 +38,13 @@ void test(T Identity, size_t WGSize, size_t NWItems) {
   queue Q;
   Q.submit([&](handler &CGH) {
     auto In = InBuf.template get_access<access::mode::read>(CGH);
+#ifdef TEST_SYCL2020_REDUCTIONS
+    auto Redu = sycl::reduction(OutBuf, CGH, Identity, BOp);
+#else
     accessor<T, Dim, access::mode::read_write, access::target::global_buffer>
         Out(OutBuf, CGH);
     auto Redu = ONEAPI::reduction(Out, Identity, BOp);
+#endif
 
     range<1> GlobalRange(NWItems);
     range<1> LocalRange(WGSize);

SYCL/Reduction/reduction_nd_s1_rw.cpp

@@ -1,9 +1,13 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
-// RUNx: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DTEST_SYCL2020_REDUCTIONS %s -o %t2020.out
+// RUN: %CPU_RUN_PLACEHOLDER %t2020.out
+// RUN: %GPU_RUN_PLACEHOLDER %t2020.out
+// RUN: %ACC_RUN_PLACEHOLDER %t2020.out
+
 // This test performs basic checks of parallel_for(nd_range, reduction, func)
 // with reductions initialized with 1-dimensional read_write accessor
 // accessing 1 element buffer.

SYCL/Reduction/reduction_placeholder.cpp

@@ -16,22 +16,25 @@
 
 using namespace cl::sycl;
 
+template <typename... Ts> class KNGroup;
+
 template <typename SpecializationKernelName, typename T, int Dim,
-          class BinaryOperation>
-void test(T Identity, size_t WGSize, size_t NWItems) {
+          class BinaryOperation, access::mode Mode>
+void testOneCase(T Identity, T Init, size_t WGSize, size_t NWItems) {
   // Initialize.
   T CorrectOut;
   BinaryOperation BOp;
 
   buffer<T, 1> OutBuf(1);
   buffer<T, 1> InBuf(NWItems);
   initInputData(InBuf, CorrectOut, Identity, BOp, NWItems);
+  if (Mode == access::mode::read_write)
+    CorrectOut = BOp(CorrectOut, Init);
 
-  (OutBuf.template get_access<access::mode::write>())[0] = Identity;
+  (OutBuf.template get_access<access::mode::write>())[0] = Init;
 
-  auto Out =
-      accessor<T, Dim, access::mode::read_write, access::target::global_buffer,
-               access::placeholder::true_t>(OutBuf);
+  auto Out = accessor<T, Dim, Mode, access::target::global_buffer,
+                      access::placeholder::true_t>(OutBuf);
   // Compute.
   queue Q;
   Q.submit([&](handler &CGH) {
@@ -58,28 +61,34 @@ void test(T Identity, size_t WGSize, size_t NWItems) {
   }
 }
 
+template <typename KernelName, typename T, int Dim, class BinaryOperation>
+void test(T Identity, T Init, size_t WGSize, size_t NWItems) {
+  testOneCase<KNGroup<KernelName, class RWCase>, T, Dim, BinaryOperation,
+              access::mode::read_write>(Identity, Init, WGSize, NWItems);
+  testOneCase<KNGroup<KernelName, class DWCase>, T, Dim, BinaryOperation,
+              access::mode::discard_write>(Identity, Init, WGSize, NWItems);
+}
+
 int main() {
   // fast atomics and fast reduce
-  test<class KernelName_LpplYqDbNN, int, 1, ONEAPI::plus<int>>(0, 49, 49 * 5);
-  test<class KernelName_FlDEESAfXYXiBZhnEDQ, int, 0, ONEAPI::plus<int>>(0, 8,
-                                                                        8);
+  test<class AtomicReduce1, int, 1, ONEAPI::plus<int>>(0, 77, 49, 49 * 5);
+  test<class AtomicReduce2, int, 0, ONEAPI::plus<int>>(0, -77, 8, 8);
 
   // fast atomics
-  test<class KernelName_caKErpdwXzEsGGkr, int, 0, ONEAPI::bit_or<int>>(0, 7,
-                                                                       7 * 3);
-  test<class KernelName_YbnKY, int, 1, ONEAPI::bit_or<int>>(0, 4, 128);
+  test<class Atomic1, int, 0, ONEAPI::bit_or<int>>(0, 233, 7, 7 * 3);
+  test<class Atomic2, int, 1, ONEAPI::bit_or<int>>(0, 177, 4, 128);
 
   // fast reduce
-  test<class KernelName_JQuAndqdovQbAHmVLm, float, 1, ONEAPI::minimum<float>>(
-      getMaximumFPValue<float>(), 5, 5 * 7);
-  test<class KernelName_MBbbTWwSc, float, 0, ONEAPI::maximum<float>>(
-      getMinimumFPValue<float>(), 4, 128);
+  test<class Reduce1, float, 1, ONEAPI::minimum<float>>(
+      getMaximumFPValue<float>(), -5.0, 5, 5 * 7);
+  test<class Reduce2, float, 0, ONEAPI::maximum<float>>(
+      getMinimumFPValue<float>(), -5.0, 4, 128);
 
   // generic algorithm
-  test<class KernelName_WpNdTbTtYt, int, 0, std::multiplies<int>>(1, 7, 7 * 5);
-  test<class KernelName_yAwH, int, 1, std::multiplies<int>>(1, 8, 16);
-  test<class KernelName_BNuHxeewzfXATi, CustomVec<short>, 0,
-       CustomVecPlus<short>>(CustomVec<short>(0), 8, 8 * 3);
+  test<class Generic1, int, 0, std::multiplies<int>>(1, 2, 7, 7 * 5);
+  test<class Generic2, int, 1, std::multiplies<int>>(1, 3, 8, 16);
+  test<class Generic13, CustomVec<short>, 0, CustomVecPlus<short>>(
+      CustomVec<short>(0), CustomVec<short>(4), 8, 8 * 3);
 
   std::cout << "Test passed\n";
   return 0;

SYCL/Reduction/reduction_queue_parallel_for.cpp

@@ -3,7 +3,13 @@
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DTEST_SYCL2020_REDUCTIONS %s -o %t2020.out
+// RUN: %CPU_RUN_PLACEHOLDER %t2020.out
+// RUN: %GPU_RUN_PLACEHOLDER %t2020.out
+// RUN: %ACC_RUN_PLACEHOLDER %t2020.out
+
 // RUNx: %HOST_RUN_PLACEHOLDER %t.out
+// RUNx: %HOST_RUN_PLACEHOLDER %t2020.out
 // TODO: Enable the test for HOST when it supports ONEAPI::reduce() and
 // barrier()
 
@@ -25,10 +31,17 @@ int main() {
   int *Sum = malloc_shared<int>(1, Q);
   *Sum = 0;
 
+#ifdef TEST_SYCL2020_REDUCTIONS
+  Q.parallel_for<class XYZ>(
+       nd_range<1>{NElems, WGSize}, sycl::reduction(Sum, std::plus<>()),
+       [=](nd_item<1> It, auto &Sum) { Sum += Data[It.get_global_id(0)]; })
+      .wait();
+#else
   Q.parallel_for<class XYZ>(
        nd_range<1>{NElems, WGSize}, ONEAPI::reduction(Sum, ONEAPI::plus<>()),
        [=](nd_item<1> It, auto &Sum) { Sum += Data[It.get_global_id(0)]; })
       .wait();
+#endif
 
   int ExpectedSum = (NElems - 1) * NElems / 2;
   int Error = 0;