diff --git a/SYCL/Reduction/reduction_big_data.cpp b/SYCL/Reduction/reduction_big_data.cpp
index 1c8d67068b..e7fd13d1a5 100644
--- a/SYCL/Reduction/reduction_big_data.cpp
+++ b/SYCL/Reduction/reduction_big_data.cpp
@@ -4,7 +4,7 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 
 // RUNx: %HOST_RUN_PLACEHOLDER %t.out
-// TODO: Enable the test for HOST when it supports ONEAPI::reduce() and
+// TODO: Enable the test for HOST when it supports ext::oneapi::reduce() and
 // barrier()
 
 // This test performs basic checks of parallel_for(nd_range, reduction, func)
@@ -12,9 +12,8 @@
 // errors.
 
 #include "reduction_utils.hpp"
-#include <CL/sycl.hpp>
+
 #include <algorithm>
-#include <cassert>
 
 using namespace cl::sycl;
 
@@ -27,14 +26,14 @@ size_t getSafeMaxWGSize(size_t MaxWGSize, size_t MemSize, size_t OneElemSize) {
   return std::min(MaxNumElems / 2, MaxWGSize);
 }
 
-template <typename KernelName, typename T, int Dim, class BinaryOperation>
-void test(queue &Q, T Identity) {
-  device Device = Q.get_device();
+template <typename Name, typename T, int Dim, class BinaryOperation>
+int test(queue &Q, T Identity) {
+  // It seems enough to test just one case - SYCL2020 reduction.
+  constexpr bool IsSYCL2020 = true;
 
+  device Device = Q.get_device();
   std::size_t MaxWGSize = Device.get_info<info::device::max_work_group_size>();
   std::size_t LocalMemSize = Device.get_info<info::device::local_mem_size>();
-  std::cout << "Detected device::max_work_group_size = " << MaxWGSize << "\n";
-  std::cout << "Detected device::local_mem_size = " << LocalMemSize << "\n";
 
   size_t WGSize = getSafeMaxWGSize(MaxWGSize, LocalMemSize, sizeof(T));
 
@@ -44,15 +43,12 @@ void test(queue &Q, T Identity) {
   size_t NWorkItems = std::min(WGSize * MaxWGSize + 1, MaxGlobalRange);
 
   size_t NWorkGroups = (NWorkItems - 1) / WGSize + 1;
-  range<1> GlobalRange(NWorkGroups * WGSize);
-  range<1> LocalRange(WGSize);
-  nd_range<1> NDRange(GlobalRange, LocalRange);
-  std::cout << "Running the test with: GlobalRange = " << (NWorkGroups * WGSize)
-            << ", LocalRange = " << WGSize << ", NWorkItems = " << NWorkItems
-            << "\n";
+  nd_range<1> NDRange(range<1>{NWorkGroups * WGSize}, range<1>{WGSize});
+  printTestLabel<T, BinaryOperation>(IsSYCL2020, NDRange);
 
   buffer<T, 1> InBuf(NWorkItems);
   buffer<T, 1> OutBuf(1);
+  (OutBuf.template get_access<access::mode::write>())[0] = Identity;
 
   // Initialize.
   BinaryOperation BOp;
@@ -62,25 +58,17 @@ void test(queue &Q, T Identity) {
   // Compute.
   Q.submit([&](handler &CGH) {
     auto In = InBuf.template get_access<access::mode::read>(CGH);
-    accessor<T, Dim, access::mode::discard_write, access::target::global_buffer>
-        Out(OutBuf, CGH);
-    CGH.parallel_for<KernelName>(NDRange, ONEAPI::reduction(Out, Identity, BOp),
-                                 [=](nd_item<1> NDIt, auto &Sum) {
-                                   if (NDIt.get_global_linear_id() < NWorkItems)
-                                     Sum.combine(
-                                         In[NDIt.get_global_linear_id()]);
-                                 });
+    CGH.parallel_for<Name>(NDRange, sycl::reduction(OutBuf, CGH, Identity, BOp),
+                           [=](nd_item<1> NDIt, auto &Sum) {
+                             if (NDIt.get_global_linear_id() < NWorkItems)
+                               Sum.combine(In[NDIt.get_global_linear_id()]);
+                           });
   });
 
   // Check correctness.
   auto Out = OutBuf.template get_access<access::mode::read>();
   T ComputedOut = *(Out.get_pointer());
-  if (ComputedOut != CorrectOut) {
-    std::cout << "Computed value: " << ComputedOut
-              << ", Expected value: " << CorrectOut << "\n";
-    assert(0 && "Wrong value.");
-  }
-  std::cout << "Test case passed\n\n";
+  return checkResults(Q, IsSYCL2020, BOp, NDRange, ComputedOut, CorrectOut);
 }
 
 template <typename T> struct BigCustomVec : public CustomVec<T> {
@@ -99,11 +87,13 @@ template <class T> struct BigCustomVecPlus {
 
 int main() {
   queue Q;
-  test<class Test1, float, 0, ONEAPI::maximum<>>(Q, getMinimumFPValue<float>());
+  printDeviceInfo(Q);
+  int NumErrors = test<class A1, float, 0, ext::oneapi::maximum<>>(
+      Q, getMinimumFPValue<float>());
 
   using BCV = BigCustomVec<long long>;
-  test<class Test2, BCV, 1, BigCustomVecPlus<long long>>(Q, BCV(0));
+  NumErrors += test<class A2, BCV, 1, BigCustomVecPlus<long long>>(Q, BCV(0));
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_nd_N_vars.cpp b/SYCL/Reduction/reduction_nd_N_vars.cpp
index 1fc8cfadc7..0f4b4ead79 100644
--- a/SYCL/Reduction/reduction_nd_N_vars.cpp
+++ b/SYCL/Reduction/reduction_nd_N_vars.cpp
@@ -4,23 +4,13 @@
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
 // TODO: The test irregularly reports incorrect.
-// UNSUPPORTED: TEMPORARY_DISABLED
+// REQUIRES: TEMPORARY_DISABLED
 
 // This test checks handling of parallel_for() accepting nd_range and
 // two or more reductions.
 
 #include "reduction_utils.hpp"
 
-#include <CL/sycl.hpp>
-
-#include <cassert>
-#include <cmath>
-#include <cstdint>
-#include <cstdio>
-#include <cstdlib>
-#include <numeric>
-#include <string>
-
 using namespace cl::sycl;
 
 template <typename... Ts> class KNameGroup;
@@ -29,22 +19,15 @@ template <typename T, bool B> class KName;
 constexpr access::mode RW = access::mode::read_write;
 constexpr access::mode DW = access::mode::discard_write;
 
-template <typename T>
-bool cherkResultIsExpected(int TestCaseNum, T Expected, T Computed,
-                           bool IsSYCL2020) {
-  bool Success;
-  if (!std::is_floating_point<T>::value)
-    Success = (Expected == Computed);
-  else
-    Success = std::abs((Expected / Computed) - 1) < 0.5;
-
-  if (!Success) {
-    std::cerr << "Is SYCL2020 mode: " << IsSYCL2020 << std::endl;
-    std::cerr << TestCaseNum << ": Expected value = " << Expected
-              << ", Computed value = " << Computed << "\n";
-  }
-
-  return Success;
+template <typename RangeT>
+void printNVarsTestLabel(bool IsSYCL2020, const RangeT &Range,
+                         bool ToCERR = false) {
+  std::ostream &OS = ToCERR ? std::cerr : std::cout;
+  std::string Mode = IsSYCL2020 ? "SYCL2020" : "ONEAPI  ";
+  OS << (ToCERR ? "Error" : "Start") << ": Mode=" << Mode
+     << ", Range=" << Range;
+  if (!ToCERR)
+    OS << std::endl;
 }
 
 // Returns 0 if the test case passed. Otherwise, some non-zero value.
@@ -58,6 +41,10 @@ int testOne(queue &Q, T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
             T3 IdentityVal3, T3 InitVal3, BinaryOperation3 BOp3,
             T4 IdentityVal4, T3 InitVal4, BinaryOperation4 BOp4,
             usm::alloc AllocType4, size_t NWorkItems, size_t WGSize) {
+
+  auto NDR = nd_range<1>{range<1>(NWorkItems), range<1>{WGSize}};
+  printNVarsTestLabel<>(IsSYCL2020, NDR);
+
   buffer<T1, 1> InBuf1(NWorkItems);
   buffer<T2, 1> InBuf2(NWorkItems);
   buffer<T3, 1> InBuf3(NWorkItems);
@@ -120,64 +107,33 @@ int testOne(queue &Q, T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
     }
   }
 
-  auto NDR = nd_range<1>{range<1>(NWorkItems), range<1>{WGSize}};
-  if constexpr (IsSYCL2020) {
-    Q.submit([&](handler &CGH) {
-       auto In1 = InBuf1.template get_access<access::mode::read>(CGH);
-       auto In2 = InBuf2.template get_access<access::mode::read>(CGH);
-       auto In3 = InBuf3.template get_access<access::mode::read>(CGH);
-       auto In4 = InBuf4.template get_access<access::mode::read>(CGH);
-
-       auto Redu1 = sycl::reduction(OutBuf1, CGH, IdentityVal1, BOp1,
-                                    getPropertyList<Mode1>());
-       auto Redu2 = sycl::reduction(OutBuf2, CGH, IdentityVal2, BOp2,
-                                    getPropertyList<Mode2>());
-       auto Redu3 = sycl::reduction(OutBuf3, CGH, IdentityVal3, BOp3,
-                                    getPropertyList<Mode3>());
-       auto Redu4 =
-           sycl::reduction(Out4, IdentityVal4, BOp4, getPropertyList<Mode4>());
-
-       auto Lambda = [=](nd_item<1> NDIt, auto &Sum1, auto &Sum2, auto &Sum3,
-                         auto &Sum4) {
-         size_t I = NDIt.get_global_id(0);
-         Sum1.combine(In1[I]);
-         Sum2.combine(In2[I]);
-         Sum3.combine(In3[I]);
-         Sum4.combine(In4[I]);
-       };
-       CGH.parallel_for<Name>(NDR, Redu1, Redu2, Redu3, Redu4, Lambda);
-     }).wait();
-  } else {
-    // Test ONEAPI reductions
-    Q.submit([&](handler &CGH) {
-       auto In1 = InBuf1.template get_access<access::mode::read>(CGH);
-       auto In2 = InBuf2.template get_access<access::mode::read>(CGH);
-       auto In3 = InBuf3.template get_access<access::mode::read>(CGH);
-       auto In4 = InBuf4.template get_access<access::mode::read>(CGH);
-
-       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);
-
-       auto Lambda = [=](nd_item<1> NDIt, auto &Sum1, auto &Sum2, auto &Sum3,
-                         auto &Sum4) {
-         size_t I = NDIt.get_global_id(0);
-         Sum1.combine(In1[I]);
-         Sum2.combine(In2[I]);
-         Sum3.combine(In3[I]);
-         Sum4.combine(In4[I]);
-       };
-       CGH.parallel_for<Name>(NDR, Redu1, Redu2, Redu3, Redu4, Lambda);
-     }).wait();
-  }
+  Q.submit([&](handler &CGH) {
+     auto In1 = InBuf1.template get_access<access::mode::read>(CGH);
+     auto In2 = InBuf2.template get_access<access::mode::read>(CGH);
+     auto In3 = InBuf3.template get_access<access::mode::read>(CGH);
+     auto In4 = InBuf4.template get_access<access::mode::read>(CGH);
+
+     auto Redu1 =
+         createReduction<IsSYCL2020, Mode1>(OutBuf1, CGH, IdentityVal1, BOp1);
+     auto Redu2 =
+         createReduction<IsSYCL2020, Mode2>(OutBuf2, CGH, IdentityVal2, BOp2);
+     auto Redu3 =
+         createReduction<IsSYCL2020, Mode3>(OutBuf3, CGH, IdentityVal3, BOp3);
+     auto Redu4 = createReduction<IsSYCL2020, Mode4>(Out4, IdentityVal4, BOp4);
+
+     auto Lambda = [=](nd_item<1> NDIt, auto &Sum1, auto &Sum2, auto &Sum3,
+                       auto &Sum4) {
+       size_t I = NDIt.get_global_id(0);
+       Sum1.combine(In1[I]);
+       Sum2.combine(In2[I]);
+       Sum3.combine(In3[I]);
+       Sum4.combine(In4[I]);
+     };
+     CGH.parallel_for<Name>(NDR, Redu1, Redu2, Redu3, Redu4, Lambda);
+   }).wait();
 
   // Check the results and free memory.
-  int Error = 0;
+  int NumErrors = 0;
   {
     auto Out1 = OutBuf1.template get_access<access::mode::read>();
     auto Out2 = OutBuf2.template get_access<access::mode::read>();
@@ -195,22 +151,23 @@ int testOne(queue &Q, T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
       Out4Val = *Out4;
     }
 
-    Error += cherkResultIsExpected(1, CorrectOut1, Out1[0], IsSYCL2020) ? 0 : 1;
-    Error += cherkResultIsExpected(2, CorrectOut2, Out2[0], IsSYCL2020) ? 0 : 1;
-    Error += cherkResultIsExpected(3, CorrectOut3, Out3[0], IsSYCL2020) ? 0 : 1;
-    Error += cherkResultIsExpected(4, CorrectOut4, Out4Val, IsSYCL2020) ? 0 : 1;
+    std::string AddInfo = "TestCase=";
+    NumErrors += checkResults(Q, IsSYCL2020, BOp1, NDR, Out1[0], CorrectOut1,
+                              AddInfo + std::to_string(1));
+    NumErrors += checkResults(Q, IsSYCL2020, BOp2, NDR, Out2[0], CorrectOut2,
+                              AddInfo + std::to_string(2));
+    NumErrors += checkResults(Q, IsSYCL2020, BOp3, NDR, Out3[0], CorrectOut3,
+                              AddInfo + std::to_string(3));
+    NumErrors += checkResults(Q, IsSYCL2020, BOp4, NDR, Out4Val, CorrectOut4,
+                              AddInfo + std::to_string(4));
     free(Out4, Q.get_context());
   }
 
-  if (Error)
-    std::cerr << "The test failed for nd_range(" << NWorkItems << "," << WGSize
-              << ")\n\n";
-
-  return Error;
+  return NumErrors;
 }
 
 // Tests both implementations of reduction:
-// sycl::reduction and sycl::ONEAPI::reduction
+// sycl::reduction and sycl::ext::oneapi::reduction
 template <class Name, typename T1, access::mode Mode1, typename T2,
           access::mode Mode2, typename T3, access::mode Mode3, typename T4,
           access::mode Mode4, class BinaryOperation1, class BinaryOperation2,
@@ -236,18 +193,16 @@ int testBoth(queue &Q, T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
 
 int main() {
   queue Q;
+  printDeviceInfo(Q);
   int Error = testBoth<class Case1, float, DW, int, RW, short, RW, int, RW>(
-      Q, 0, 1000, std::plus<float>{}, 0, 2000, std::plus<>{}, 0, 4000,
+      Q, 0, 1000, std::plus<>{}, 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 += testBoth<class Case2, float, RW, int, RW, short, DW, int, DW>(
-      Q, 0, 1000, std::plus<float>{}, 0, 2000, std::plus<>{}, 0, 4000, Add, 0,
-      8000, std::plus<>{}, usm::alloc::device, 5 * (256 + 1), 5);
+      Q, 0, 1000, std::plus<>{}, 0, 2000, std::plus<>{}, 0, 4000, Add, 0, 8000,
+      std::plus<>{}, usm::alloc::device, 5 * (256 + 1), 5);
 
-  if (!Error)
-    std::cout << "Test passed\n";
-  else
-    std::cerr << Error << " test-cases failed\n";
+  printFinalStatus(Error);
   return Error;
 }
diff --git a/SYCL/Reduction/reduction_nd_conditional.cpp b/SYCL/Reduction/reduction_nd_conditional.cpp
index 48c2cf56f3..d1bf4c0f05 100644
--- a/SYCL/Reduction/reduction_nd_conditional.cpp
+++ b/SYCL/Reduction/reduction_nd_conditional.cpp
@@ -7,8 +7,7 @@
 // This test performs basic checks of parallel_for(nd_range, reduction, func)
 // with reduction and conditional increment of the reduction variable.
 
-#include <CL/sycl.hpp>
-#include <cassert>
+#include "reduction_utils.hpp"
 
 using namespace cl::sycl;
 
@@ -18,10 +17,7 @@ void initInputData(buffer<T, 1> &InBuf, T &ExpectedOut, T Identity,
   ExpectedOut = Identity;
   auto In = InBuf.template get_access<access::mode::write>();
   for (int I = 0; I < N; ++I) {
-    if (std::is_same<BinaryOperation, std::multiplies<T>>::value)
-      In[I] = 1 + (((I % 37) == 0) ? 1 : 0);
-    else
-      In[I] = I + 1 + 1.1;
+    In[I] = I + 1;
 
     if (I < 2)
       ExpectedOut = BOp(ExpectedOut, 99);
@@ -32,79 +28,51 @@ void initInputData(buffer<T, 1> &InBuf, T &ExpectedOut, T Identity,
   }
 };
 
-template <typename T> struct Vec {
-  Vec() : X(0), Y(0) {}
-  Vec(T X, T Y) : X(X), Y(Y) {}
-  Vec(T V) : X(V), Y(V) {}
-  bool operator==(const Vec &P) const { return P.X == X && P.Y == Y; }
-  bool operator!=(const Vec &P) const { return !(*this == P); }
-  T X;
-  T Y;
-};
-template <typename T> bool operator==(const Vec<T> &A, const Vec<T> &B) {
-  return A.X == B.X && A.Y == B.Y;
-}
-template <typename T>
-std::ostream &operator<<(std::ostream &OS, const Vec<T> &P) {
-  return OS << "(" << P.X << ", " << P.Y << ")";
-}
-
-template <class T> struct VecPlus {
-  using P = Vec<T>;
-  P operator()(const P &A, const P &B) const { return P(A.X + B.X, A.Y + B.Y); }
-};
-
-template <typename SpecializationKernelName, typename T, int Dim,
-          class BinaryOperation>
-void test(queue &Q, T Identity, size_t WGSize, size_t NWItems) {
+template <typename Name, typename T, class BinaryOperation>
+int test(queue &Q, T Identity, size_t WGSize, size_t NWItems) {
   buffer<T, 1> InBuf(NWItems);
   buffer<T, 1> OutBuf(1);
 
+  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
+  printTestLabel<T, BinaryOperation>(true /*SYCL2020*/, NDRange);
+
   // Initialize.
   BinaryOperation BOp;
   T CorrectOut;
   initInputData(InBuf, CorrectOut, Identity, BOp, NWItems);
+  (OutBuf.template get_access<access::mode::write>())[0] = Identity;
 
   // Compute.
   Q.submit([&](handler &CGH) {
     auto In = InBuf.template get_access<access::mode::read>(CGH);
-    accessor<T, Dim, access::mode::discard_write, access::target::global_buffer>
-        Out(OutBuf, CGH);
-    auto Redu = ONEAPI::reduction(Out, Identity, BOp);
-
-    range<1> GlobalRange(NWItems);
-    range<1> LocalRange(WGSize);
-    nd_range<1> NDRange(GlobalRange, LocalRange);
-    CGH.parallel_for<SpecializationKernelName>(
-        NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-          size_t I = NDIt.get_global_linear_id();
-          if (I < 2)
-            Sum.combine(T(99));
-          else if (I % 3)
-            Sum.combine(In[I]);
-          else
-            ; // do nothing.
-        });
+    auto Redu = sycl::reduction(OutBuf, CGH, Identity, BOp);
+    CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
+      size_t I = NDIt.get_global_linear_id();
+      if (I < 2)
+        Sum.combine(T(99));
+      else if (I % 3)
+        Sum.combine(In[I]);
+      else
+        ; // do nothing.
+    });
   });
 
   // Check correctness.
   auto Out = OutBuf.template get_access<access::mode::read>();
   T ComputedOut = *(Out.get_pointer());
-  if (ComputedOut != CorrectOut) {
-    std::cout << "NWItems = " << NWItems << ", WGSize = " << WGSize << "\n";
-    std::cout << "Computed value: " << ComputedOut
-              << ", Expected value: " << CorrectOut << "\n";
-    assert(0 && "Wrong value.");
-  }
+  return checkResults(Q, true /*SYCL2020*/, BOp, NDRange, ComputedOut,
+                      CorrectOut);
 }
 
 int main() {
   queue Q;
-  test<class A, int, 0, ONEAPI::plus<int>>(Q, 0, 2, 2);
-  test<class B, int, 1, ONEAPI::plus<int>>(Q, 0, 7, 7);
-  test<class C, int, 0, ONEAPI::plus<int>>(Q, 0, 2, 64);
-  test<class D, short, 1, ONEAPI::plus<short>>(Q, 0, 16, 256);
+  printDeviceInfo(Q);
+
+  int NumErrors = 0;
+  NumErrors += test<class A, unsigned char, std::plus<>>(Q, 0, 2, 2);
+  NumErrors += test<class B, short, std::plus<>>(Q, 0, 7, 7);
+  NumErrors += test<class C, int, std::plus<>>(Q, 0, 1, 1025);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_nd_ext_double.cpp b/SYCL/Reduction/reduction_nd_ext_double.cpp
index aacbce14cb..88b375ca16 100644
--- a/SYCL/Reduction/reduction_nd_ext_double.cpp
+++ b/SYCL/Reduction/reduction_nd_ext_double.cpp
@@ -11,4 +11,4 @@
 
 #include "reduction_nd_ext_type.hpp"
 
-int main() { return runTests<double>("cl_khr_fp64"); }
+int main() { return runTests<double>(sycl::aspect::fp64); }
diff --git a/SYCL/Reduction/reduction_nd_ext_half.cpp b/SYCL/Reduction/reduction_nd_ext_half.cpp
index 7bb9f6f543..e0e88c05ca 100644
--- a/SYCL/Reduction/reduction_nd_ext_half.cpp
+++ b/SYCL/Reduction/reduction_nd_ext_half.cpp
@@ -16,4 +16,4 @@
 
 #include "reduction_nd_ext_type.hpp"
 
-int main() { return runTests<half>("cl_khr_fp16"); }
+int main() { return runTests<sycl::half>(sycl::aspect::fp16); }
diff --git a/SYCL/Reduction/reduction_nd_ext_type.hpp b/SYCL/Reduction/reduction_nd_ext_type.hpp
index 0ee925a615..f30768c428 100644
--- a/SYCL/Reduction/reduction_nd_ext_type.hpp
+++ b/SYCL/Reduction/reduction_nd_ext_type.hpp
@@ -2,103 +2,47 @@
 // with types that may require additional runtime checks for extensions
 // supported by the device, e.g. 'half' or 'double'
 
-#include "reduction_utils.hpp"
-#include <CL/sycl.hpp>
-#include <cassert>
+#include "reduction_nd_range_scalar.hpp"
 
 using namespace cl::sycl;
 
-template <typename T, bool B> class KName;
-constexpr access::mode RW = access::mode::read_write;
-constexpr access::mode DW = access::mode::discard_write;
+int NumErrors = 0;
 
-template <typename Name, bool IsSYCL2020Mode, typename T, int Dim,
-          access::mode Mode, class BinaryOperation>
-void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  buffer<T, 1> InBuf(NWItems);
-  buffer<T, 1> OutBuf(1);
-
-  // Initialize.
-  BinaryOperation BOp;
-  T CorrectOut;
-  initInputData(InBuf, CorrectOut, Identity, BOp, NWItems);
-  if (Mode == access::mode::read_write)
-    CorrectOut = BOp(CorrectOut, Init);
-
-  (OutBuf.template get_access<access::mode::write>())[0] = Init;
-
-  // Compute.
+template <typename Name, typename T, access::mode Mode, class BinaryOperation>
+void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, size_t WGSize,
+           size_t NWItems) {
   nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
-  if constexpr (IsSYCL2020Mode) {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      auto Redu =
-          sycl::reduction(OutBuf, CGH, Identity, BOp, getPropertyList<Mode>());
-      CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-        Sum.combine(In[NDIt.get_global_linear_id()]);
-      });
-    });
-  } else {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      accessor<T, Dim, Mode, access::target::global_buffer> Out(OutBuf, CGH);
-      auto Redu = ONEAPI::reduction(Out, Identity, BOp);
-
-      CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-        Sum.combine(In[NDIt.get_global_linear_id()]);
-      });
-    });
-  }
-
-  // Check correctness.
-  auto Out = OutBuf.template get_access<access::mode::read>();
-  T ComputedOut = *(Out.get_pointer());
-  T MaxDiff = 4 * std::numeric_limits<T>::epsilon() *
-              std::fabs(ComputedOut + CorrectOut);
-  T CompDiff = std::fabs(static_cast<T>(ComputedOut - CorrectOut));
-  if (CompDiff > MaxDiff) {
-    std::cerr << "NWItems = " << NWItems << ", WGSize = " << WGSize << "\n";
-    std::cerr << "Computed value: " << ComputedOut
-              << ", Expected value: " << CorrectOut << ", MaxDiff = " << MaxDiff
-              << ", ComputedDiff = " << CompDiff << std::endl;
-    assert(0 && "Wrong value.");
-  }
-}
-
-template <typename Name, typename T, int Dim, access::mode Mode,
-          class BinaryOperation>
-void testBoth(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  test<KName<Name, false>, false, T, Dim, Mode, BinaryOperation>(
-      Q, Identity, Init, WGSize, NWItems);
-
-  test<KName<Name, true>, true, T, Dim, Mode, BinaryOperation>(
-      Q, Identity, Init, WGSize, NWItems);
+  NumErrors += testBoth<Name, Mode>(Q, Identity, Init, BOp, NDRange);
 }
 
-template <typename T> int runTests(const string_class &ExtensionName) {
+template <typename T> int runTests(sycl::aspect ExtAspect) {
   queue Q;
+  printDeviceInfo(Q);
   device D = Q.get_device();
-  if (!D.is_host() && !D.has_extension(ExtensionName)) {
+  if (!D.is_host() && !D.has(ExtAspect)) {
     std::cout << "Test skipped\n";
     return 0;
   }
 
-  testBoth<class A, T, 1, RW, std::multiplies<T>>(Q, 1, 77, 4, 4);
+  constexpr access::mode RW = access::mode::read_write;
+  constexpr access::mode DW = access::mode::discard_write;
+
+  tests<class A1, T, DW>(Q, 1, 77, std::multiplies<T>{}, 4, 4);
+  tests<class A2, T, RW>(Q, 1, 77, std::multiplies<T>{}, 4, 8);
+
+  tests<class B1, T, DW>(Q, 0, 77, std::plus<T>{}, 4, 32);
+  tests<class B2, T, RW>(Q, 0, 33, std::plus<T>{}, 3, 3 * 5);
 
-  testBoth<class B1, T, 0, DW, ONEAPI::plus<T>>(Q, 0, 77, 4, 64);
-  testBoth<class B2, T, 1, RW, ONEAPI::plus<>>(Q, 0, 33, 3, 3 * 5);
+  tests<class C1, T, DW>(Q, getMaximumFPValue<T>(), -10.0,
+                         ext::oneapi::minimum<T>{}, 7, 7 * 512);
+  tests<class C2, T, RW>(Q, getMaximumFPValue<T>(), 99.0,
+                         ext::oneapi::minimum<T>{}, 7, 7);
 
-  testBoth<class C1, T, 0, RW, ONEAPI::minimum<T>>(Q, getMaximumFPValue<T>(),
-                                                   -10.0, 7, 7);
-  testBoth<class C2, T, 0, RW, ONEAPI::minimum<T>>(Q, getMaximumFPValue<T>(),
-                                                   99.0, 7, 7);
-  testBoth<class C3, T, 1, DW, ONEAPI::minimum<>>(Q, getMaximumFPValue<T>(),
-                                                  -99.0, 3, 3);
+  tests<class D1, T, DW>(Q, getMinimumFPValue<T>(), 99.0,
+                         ext::oneapi::maximum<>{}, 3, 3);
+  tests<class D2, T, RW>(Q, getMinimumFPValue<T>(), 99.0,
+                         ext::oneapi::maximum<>{}, 7, 7 * 5);
 
-  testBoth<class D1, T, 0, DW, ONEAPI::maximum<>>(Q, getMinimumFPValue<T>(),
-                                                  99.0, 3, 3);
-  testBoth<class D2, T, 1, RW, ONEAPI::maximum<T>>(Q, getMinimumFPValue<T>(),
-                                                   99.0, 7, 7 * 5);
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_nd_lambda.cpp b/SYCL/Reduction/reduction_nd_lambda.cpp
index 0fe7b29779..334385b65b 100644
--- a/SYCL/Reduction/reduction_nd_lambda.cpp
+++ b/SYCL/Reduction/reduction_nd_lambda.cpp
@@ -6,72 +6,42 @@
 
 // This test performs basic checks of parallel_for(nd_range, reduction, lambda)
 
-#include "reduction_utils.hpp"
-#include <CL/sycl.hpp>
-#include <cassert>
+#include "reduction_nd_range_scalar.hpp"
 
 using namespace cl::sycl;
 
-// Note that this function is created only to test that if the accessor
-// object passed to ONEAPI::reduction is destroyed right after
-// ONEAPI::reduction creation, then the reduction still works properly,
-// i.e. it holds a COPY of user's accessor.
-template <typename T, typename BOpT>
-auto createReduction(sycl::buffer<T, 1> Buffer, handler &CGH, T Identity,
-                     BOpT BOp) {
-  auto Acc = Buffer.template get_access<access::mode::discard_write>(CGH);
-  return ONEAPI::reduction(Acc, Identity, BOp);
-}
-
-template <class KernelName, typename T, class BinaryOperation>
-void test(queue &Q, T Identity, BinaryOperation BOp, size_t WGSize,
-          size_t NWItems) {
-  buffer<T, 1> InBuf(NWItems);
-  buffer<T, 1> OutBuf(1);
-
-  // Initialize.
-  T CorrectOut;
-  initInputData(InBuf, CorrectOut, Identity, BOp, NWItems);
-
-  // Compute.
-  Q.submit([&](handler &CGH) {
-    auto In = InBuf.template get_access<access::mode::read>(CGH);
-    auto Redu = createReduction(OutBuf, CGH, Identity, BOp);
+template <typename T, access::mode M> class MName;
 
-    nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
-    CGH.parallel_for<KernelName>(NDRange, Redu,
-                                 [=](nd_item<1> NDIt, auto &Sum) {
-                                   Sum.combine(In[NDIt.get_global_linear_id()]);
-                                 });
-  });
+int NumErrors = 0;
 
-  // Check correctness.
-  auto Out = OutBuf.template get_access<access::mode::read>();
-  T ComputedOut = *(Out.get_pointer());
-  if (ComputedOut != CorrectOut) {
-    std::cerr << "NWItems = " << NWItems << ", WGSize = " << WGSize << "\n";
-    std::cerr << "Computed value: " << ComputedOut
-              << ", Expected value: " << CorrectOut << "\n";
-    assert(0 && "Wrong value.");
-  }
+template <typename Name, typename T, class BinaryOperation>
+void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, size_t WGSize,
+           size_t NWItems) {
+  constexpr access::mode DW = access::mode::discard_write;
+  constexpr access::mode RW = access::mode::read_write;
+  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
+  NumErrors += testBoth<MName<Name, DW>, DW>(Q, Identity, Init, BOp, NDRange);
+  NumErrors += testBoth<MName<Name, RW>, RW>(Q, Identity, Init, BOp, NDRange);
 }
 
 int main() {
   queue Q;
-  test<class AddTestName, int>(
-      Q, 0, [](auto x, auto y) { return (x + y); }, 1, 1024);
-  test<class MulTestName, int>(
-      Q, 0, [](auto x, auto y) { return (x * y); }, 8, 32);
+  printDeviceInfo(Q);
+  tests<class A1, int>(
+      Q, 0, 9, [](auto x, auto y) { return (x + y); }, 1, 1024);
+  tests<class A2, uint64_t>(
+      Q, 1, 2, [](auto x, auto y) { return (x * y); }, 8, 16);
 
   // Check with CUSTOM type.
-  test<class CustomAddTestname, CustomVec<long long>>(
-      Q, CustomVec<long long>(0),
-      [](auto x, auto y) {
+  using CV = CustomVec<long long>;
+  tests<class A3, CV>(
+      Q, CV(0), CV(2021),
+      [](auto X, auto Y) {
         CustomVecPlus<long long> BOp;
-        return BOp(x, y);
+        return BOp(X, Y);
       },
       4, 64);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_nd_range_scalar.hpp b/SYCL/Reduction/reduction_nd_range_scalar.hpp
new file mode 100644
index 0000000000..cdfdc897dc
--- /dev/null
+++ b/SYCL/Reduction/reduction_nd_range_scalar.hpp
@@ -0,0 +1,59 @@
+// This test performs basic checks of parallel_for(range<Dims>, reduction, func)
+// with reductions initialized with 1-dimensional buffer/accessor
+// accessing a scalar holding the reduction result.
+
+#include "reduction_utils.hpp"
+
+using namespace cl::sycl;
+
+template <typename T, bool B> class KName;
+
+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 nd_range<Dims> &Range) {
+  printTestLabel<T, BinaryOperation>(IsSYCL2020, Range);
+
+  // Skip the test for such big arrays now.
+  constexpr size_t TwoGB = 2LL * 1024 * 1024 * 1024;
+  range<Dims> GlobalRange = Range.get_global_range();
+  if (GlobalRange.size() > TwoGB)
+    return 0;
+
+  buffer<T, Dims> InBuf(GlobalRange);
+  buffer<T, 1> OutBuf(1);
+
+  // Initialize.
+  T CorrectOut;
+  initInputData(InBuf, CorrectOut, Identity, BOp, GlobalRange);
+  if constexpr (Mode == access::mode::read_write) {
+    CorrectOut = BOp(CorrectOut, Init);
+  }
+
+  // The value assigned here must be discarded (if IsReadWrite is true).
+  // Verify that it is really discarded and assign some value.
+  (OutBuf.template get_access<access::mode::write>())[0] = Init;
+
+  // Compute.
+  Q.submit([&](handler &CGH) {
+    auto In = InBuf.template get_access<access::mode::read>(CGH);
+    auto Redu =
+        createReduction<IsSYCL2020, Mode, AccDim>(OutBuf, CGH, Identity, BOp);
+    CGH.parallel_for<Name>(Range, Redu, [=](nd_item<Dims> NDIt, auto &Sum) {
+      Sum.combine(In[NDIt.get_global_id()]);
+    });
+  });
+
+  // Check correctness.
+  auto Out = OutBuf.template get_access<access::mode::read>();
+  T ComputedOut = *(Out.get_pointer());
+  return checkResults(Q, IsSYCL2020, BOp, Range, ComputedOut, CorrectOut);
+}
+
+template <typename Name, access::mode Mode, typename T, class BinaryOperation,
+          int Dims>
+int testBoth(queue &Q, T Identity, T Init, BinaryOperation BOp,
+             const nd_range<Dims> &Range) {
+  return test<KName<Name, false>, false, Mode>(Q, Identity, Init, BOp, Range) +
+         test<KName<Name, true>, true, Mode>(Q, Identity, Init, BOp, Range);
+}
diff --git a/SYCL/Reduction/reduction_nd_s0_dw.cpp b/SYCL/Reduction/reduction_nd_s0_dw.cpp
index 43ce0f6566..6dda0ca243 100644
--- a/SYCL/Reduction/reduction_nd_s0_dw.cpp
+++ b/SYCL/Reduction/reduction_nd_s0_dw.cpp
@@ -6,104 +6,49 @@
 // This test performs basic checks of parallel_for(nd_range, reduction, func)
 // with reductions initialized with 0-dimensional discard_write accessor.
 
-#include "reduction_utils.hpp"
-#include <CL/sycl.hpp>
-#include <cassert>
+#include "reduction_nd_range_scalar.hpp"
 
 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 T, bool B> class KName;
-
-template <typename Name, bool IsSYCL2020Mode, typename T, class BinaryOperation>
-void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  buffer<T, 1> InBuf(NWItems);
-  buffer<T, 1, CustomAllocator> OutBuf(1);
-
-  // Initialize.
-  BinaryOperation BOp;
-  T CorrectOut;
-  initInputData(InBuf, CorrectOut, Identity, BOp, NWItems);
-
-  // This assignment is discarded, but is still needed to verify
-  // that is discarded.
-  (OutBuf.template get_access<access::mode::write>())[0] = Init;
-
-  // Compute.
-
-  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
-  if constexpr (IsSYCL2020Mode) {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      property_list PropList(property::reduction::initialize_to_identity{});
-      auto Redu = sycl::reduction(OutBuf, CGH, Identity, BOp, PropList);
-
-      CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-        Sum.combine(In[NDIt.get_global_linear_id()]);
-      });
-    });
-  } else {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      accessor<T, 0, access::mode::discard_write, access::target::global_buffer>
-          Out(OutBuf, CGH);
-      auto Redu = ONEAPI::reduction(Out, Identity, BOp);
-
-      CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-        Sum.combine(In[NDIt.get_global_linear_id()]);
-      });
-    });
-  }
-
-  // Check correctness.
-  auto Out = OutBuf.template get_access<access::mode::read>();
-  T ComputedOut = *(Out.get_pointer());
-  if (ComputedOut != CorrectOut) {
-    std::cout << "NWItems = " << NWItems << ", WGSize = " << WGSize << "\n";
-    std::cout << "Computed value: " << ComputedOut
-              << ", Expected value: " << CorrectOut << "\n";
-    assert(0 && "Wrong value.");
-  }
-}
+int NumErrors = 0;
 
 template <typename Name, typename T, class BinaryOperation>
-void testBoth(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  test<KName<Name, false>, false, T, BinaryOperation>(Q, Identity, Init, WGSize,
-                                                      NWItems);
-  test<KName<Name, true>, true, T, BinaryOperation>(Q, Identity, Init, WGSize,
-                                                    NWItems);
+void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, size_t WGSize,
+           size_t NWItems) {
+  constexpr access::mode DW = access::mode::discard_write;
+  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
+  NumErrors +=
+      test<Name, false /*SYCL2020*/, DW, 0>(Q, Identity, Init, BOp, NDRange);
 }
 
 int main() {
   queue Q;
+  printDeviceInfo(Q);
 
   // Check some non power-of-two work-group sizes.
-  testBoth<class A1, int, ONEAPI::plus<int>>(Q, 0, 99, 1, 7);
-  testBoth<class A2, int, ONEAPI::plus<int>>(Q, 0, 99, 49, 49 * 5);
+  tests<class A1, int>(Q, 0, 99, std::plus<int>{}, 1, 7);
+  tests<class A2, int>(Q, 0, 99, std::plus<int>{}, 49, 49 * 5);
 
   // Try some power-of-two work-group sizes.
-  testBoth<class B1, int, ONEAPI::plus<int>>(Q, 0, 99, 1, 32);
-  testBoth<class B2, int, std::multiplies<int>>(Q, 1, 99, 4, 32);
-  testBoth<class B4, int, ONEAPI::bit_xor<int>>(Q, 0, 99, 16, 256);
-  testBoth<class B5, int, ONEAPI::bit_and<int>>(Q, ~0, 99, 32, 256);
-  testBoth<class B6, int, ONEAPI::minimum<int>>(
-      Q, (std::numeric_limits<int>::max)(), -99, 64, 256);
-  testBoth<class B7, int, ONEAPI::maximum<int>>(
-      Q, (std::numeric_limits<int>::min)(), 99, 128, 256);
-  testBoth<class B8, int, ONEAPI::plus<>>(Q, 0, 99, 256, 256);
+  tests<class B1, int>(Q, 0, 99, std::plus<>{}, 1, 32);
+  tests<class B2, int>(Q, 1, 99, std::multiplies<>{}, 4, 32);
+  tests<class B4, int>(Q, 0, 99, std::bit_xor<>{}, 16, 256);
+  tests<class B5, int>(Q, ~0, 99, std::bit_and<>{}, 32, 256);
+  tests<class B6, int>(Q, (std::numeric_limits<int>::max)(), -99,
+                       ext::oneapi::minimum<>{}, 64, 256);
+  tests<class B7, int>(Q, (std::numeric_limits<int>::min)(), 99,
+                       ext::oneapi::maximum<>{}, 128, 256);
+  tests<class B8, int>(Q, 0, 99, std::plus<>{}, 256, 256);
 
   // Check with various types.
-  testBoth<class C1, float, std::multiplies<>>(Q, 1, 99, 8, 256);
-  testBoth<class C2, short, ONEAPI::minimum<>>(Q, 0x7fff, -99, 8, 256);
-  testBoth<class C3, unsigned char, ONEAPI::maximum<>>(Q, 0, 99, 8, 256);
+  tests<class C1, float>(Q, 1, 99, std::multiplies<>{}, 8, 16);
+  tests<class C2, short>(Q, 0x7fff, -99, ext::oneapi::minimum<>{}, 8, 256);
+  tests<class C3, unsigned char>(Q, 0, 99, ext::oneapi::maximum<>{}, 8, 256);
 
   // Check with CUSTOM type.
-  testBoth<class D1, CustomVec<long long>, CustomVecPlus<long long>>(
-      Q, CustomVec<long long>(0), CustomVec<long long>(99), 8, 256);
+  using CV = CustomVec<long long>;
+  tests<class D1, CV>(Q, CV(0), CV(99), CustomVecPlus<long long>{}, 8, 256);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_nd_s0_rw.cpp b/SYCL/Reduction/reduction_nd_s0_rw.cpp
index a5983f019f..a8b75b06a1 100644
--- a/SYCL/Reduction/reduction_nd_s0_rw.cpp
+++ b/SYCL/Reduction/reduction_nd_s0_rw.cpp
@@ -6,104 +6,50 @@
 // This test performs basic checks of parallel_for(nd_range, reduction, func)
 // with reductions initialized with 0-dimensional read_write accessor.
 
-#include "reduction_utils.hpp"
-#include <CL/sycl.hpp>
-#include <cassert>
+#include "reduction_nd_range_scalar.hpp"
 
 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 T, bool B> class KName;
-
-template <typename Name, bool IsSYCL2020Mode, typename T, class BinaryOperation>
-void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  buffer<T, 1> InBuf(NWItems);
-  buffer<T, 1, CustomAllocator> OutBuf(1);
-
-  // Initialize.
-  BinaryOperation BOp;
-  T CorrectOut;
-  initInputData(InBuf, CorrectOut, Identity, BOp, NWItems);
-
-  // The final reduction sum after running parallel_for() must include
-  // the original value it was initialized with before the parallel_for().
-  CorrectOut = BOp(CorrectOut, Init);
-  (OutBuf.template get_access<access::mode::write>())[0] = Init;
-
-  // Compute.
-  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
-  if constexpr (IsSYCL2020Mode) {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      auto Redu = sycl::reduction(OutBuf, CGH, Identity, BOp);
-
-      CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-        Sum.combine(In[NDIt.get_global_linear_id()]);
-      });
-    });
-  } else {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      accessor<T, 0, access::mode::read_write, access::target::global_buffer>
-          Out(OutBuf, CGH);
-      auto Redu = ONEAPI::reduction(Out, Identity, BOp);
-
-      CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-        Sum.combine(In[NDIt.get_global_linear_id()]);
-      });
-    });
-  }
-
-  // Check correctness.
-  auto Out = OutBuf.template get_access<access::mode::read>();
-  T ComputedOut = *(Out.get_pointer());
-  if (ComputedOut != CorrectOut) {
-    std::cout << "NWItems = " << NWItems << ", WGSize = " << WGSize << "\n";
-    std::cout << "Computed value: " << ComputedOut
-              << ", Expected value: " << CorrectOut << "\n";
-    assert(0 && "Wrong value.");
-  }
-}
+int NumErrors = 0;
 
 template <typename Name, typename T, class BinaryOperation>
-void testBoth(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  test<KName<Name, false>, false, T, BinaryOperation>(Q, Identity, Init, WGSize,
-                                                      NWItems);
-  test<KName<Name, true>, true, T, BinaryOperation>(Q, Identity, Init, WGSize,
-                                                    NWItems);
+void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, size_t WGSize,
+           size_t NWItems) {
+  constexpr access::mode RW = access::mode::read_write;
+  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
+  NumErrors +=
+      test<Name, false /*SYCL2020*/, RW, 0>(Q, Identity, Init, BOp, NDRange);
 }
 
 int main() {
   queue Q;
+  printDeviceInfo(Q);
 
   // Check non power-of-two work-group sizes.
-  testBoth<class A1, int, ONEAPI::plus<int>>(Q, 0, 99, 1, 7);
-  testBoth<class A2, int, ONEAPI::plus<int>>(Q, 0, -99, 49, 49 * 5);
+  tests<class A1, int>(Q, 0, 99, std::plus<int>{}, 1, 7);
+  tests<class A2, int>(Q, 0, -99, std::plus<int>{}, 49, 49 * 5);
 
   // Try some power-of-two work-group sizes.
-  testBoth<class B1, int, ONEAPI::plus<>>(Q, 0, 99, 2, 32);
-  testBoth<class B2, int, ONEAPI::plus<>>(Q, 0, 199, 32, 32);
-  testBoth<class B3, int, ONEAPI::plus<>>(Q, 0, 299, 128, 256);
-  testBoth<class B4, int, ONEAPI::plus<>>(Q, 0, 399, 256, 256);
+  tests<class B1, int>(Q, 0, 99, std::plus<>{}, 2, 32);
+  tests<class B2, int>(Q, 0, 199, std::plus<>{}, 32, 32);
+  tests<class B3, int>(Q, 0, 299, std::plus<>{}, 128, 256);
+  tests<class B4, int>(Q, 0, 399, std::plus<>{}, 256, 256);
 
   // Check with various operations and types.
-  testBoth<class C1, int, std::multiplies<int>>(Q, 1, 2, 8, 256);
-  testBoth<class C2, float, std::multiplies<float>>(Q, 1, 1.2, 8, 32);
-  testBoth<class C3, short, ONEAPI::bit_or<>>(Q, 0, 0x3400, 4, 32);
-  testBoth<class C4, int, ONEAPI::bit_xor<int>>(Q, 0, 0x12340000, 4, 32);
-  testBoth<class C5, char, ONEAPI::bit_and<>>(Q, ~0, ~0, 4, 16);
-  testBoth<class C6, int, ONEAPI::minimum<int>>(
-      Q, (std::numeric_limits<int>::max)(), -99, 8, 256);
-  testBoth<class C7, int, ONEAPI::maximum<float>>(
-      Q, (std::numeric_limits<int>::min)(), 99, 8, 256);
+  tests<class C1, int>(Q, 1, 2, std::multiplies<>{}, 8, 8);
+  tests<class C2, float>(Q, 1, 1.2, std::multiplies<>{}, 8, 16);
+  tests<class C3, short>(Q, 0, 0x3400, std::bit_or<>{}, 4, 32);
+  tests<class C4, int>(Q, 0, 0x12340000, std::bit_xor<>{}, 4, 32);
+  tests<class C5, char>(Q, ~0, ~0, std::bit_and<>{}, 4, 16);
+  tests<class C6, int>(Q, (std::numeric_limits<int>::max)(), -99,
+                       ext::oneapi::minimum<>{}, 8, 256);
+  tests<class C7, int>(Q, (std::numeric_limits<int>::min)(), 99,
+                       ext::oneapi::maximum<>{}, 8, 256);
 
   // Check with CUSTOM type.
-  testBoth<class D1, CustomVec<long long>, CustomVecPlus<long long>>(
-      Q, CustomVec<long long>(0), CustomVec<long long>(199), 8, 256);
+  using CV = CustomVec<long long>;
+  tests<class D1, CV>(Q, CV(0), CV(99), CustomVecPlus<long long>{}, 8, 256);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_nd_s1_dw.cpp b/SYCL/Reduction/reduction_nd_s1_dw.cpp
index 8745d054e3..3afbfb141b 100644
--- a/SYCL/Reduction/reduction_nd_s1_dw.cpp
+++ b/SYCL/Reduction/reduction_nd_s1_dw.cpp
@@ -8,100 +8,48 @@
 // with reductions initialized with 1-dimensional discard_write accessor
 // accessing 1 element buffer.
 
-#include "reduction_utils.hpp"
-#include <CL/sycl.hpp>
-#include <cassert>
+#include "reduction_nd_range_scalar.hpp"
 
 using namespace cl::sycl;
 
-template <typename T, bool B> class KName;
-
-template <typename Name, bool IsSYCL2020Mode, typename T, class BinaryOperation>
-void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  buffer<T, 1> InBuf(NWItems);
-  buffer<T, 1> OutBuf(1);
-
-  // Initialize.
-  BinaryOperation BOp;
-  T CorrectOut;
-  initInputData(InBuf, CorrectOut, Identity, BOp, NWItems);
-
-  // This assignment is discarded, but is still needed to verify
-  // that is discarded.
-  (OutBuf.template get_access<access::mode::write>())[0] = Init;
-
-  // Compute.
-  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
-  if constexpr (IsSYCL2020Mode) {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      property_list PropList(property::reduction::initialize_to_identity{});
-      auto Redu = sycl::reduction(OutBuf, CGH, Identity, BOp, PropList);
-
-      CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-        Sum.combine(In[NDIt.get_global_linear_id()]);
-      });
-    });
-  } else {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      accessor<T, 1, access::mode::discard_write, access::target::global_buffer>
-          Out(OutBuf, CGH);
-      auto Redu = ONEAPI::reduction(Out, Identity, BOp);
-
-      CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-        Sum.combine(In[NDIt.get_global_linear_id()]);
-      });
-    });
-  }
-
-  // Check correctness.
-  auto Out = OutBuf.template get_access<access::mode::read>();
-  T ComputedOut = *(Out.get_pointer());
-  if (ComputedOut != CorrectOut) {
-    std::cout << "NWItems = " << NWItems << ", WGSize = " << WGSize << "\n";
-    std::cout << "Computed value: " << ComputedOut
-              << ", Expected value: " << CorrectOut << "\n";
-    assert(0 && "Wrong value.");
-  }
-}
+int NumErrors = 0;
 
 template <typename Name, typename T, class BinaryOperation>
-void testBoth(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  test<KName<Name, false>, false, T, BinaryOperation>(Q, Identity, Init, WGSize,
-                                                      NWItems);
-  test<KName<Name, true>, true, T, BinaryOperation>(Q, Identity, Init, WGSize,
-                                                    NWItems);
+void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, size_t WGSize,
+           size_t NWItems) {
+  constexpr access::mode DW = access::mode::discard_write;
+  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
+  NumErrors += testBoth<Name, DW>(Q, Identity, Init, BOp, NDRange);
 }
 
 int main() {
   queue Q;
+  printDeviceInfo(Q);
 
   // Check some non power-of-two work-group sizes.
-  testBoth<class A1, int, ONEAPI::plus<int>>(Q, 0, 99, 1, 7);
-  testBoth<class A2, int, ONEAPI::plus<int>>(Q, 0, 99, 49, 49 * 5);
+  tests<class A1, int>(Q, 0, 99, std::plus<int>{}, 1, 7);
+  tests<class A2, int>(Q, 0, 99, std::plus<int>{}, 49, 49 * 5);
 
   // Try some power-of-two work-group sizes.
-  testBoth<class B1, int, ONEAPI::plus<int>>(Q, 0, 99, 1, 32);
-  testBoth<class B2, int, std::multiplies<int>>(Q, 1, 99, 4, 64);
-  testBoth<class B3, int, ONEAPI::bit_or<int>>(Q, 0, 99, 8, 128);
-  testBoth<class B4, int, ONEAPI::bit_xor<int>>(Q, 0, 99, 16, 256);
-  testBoth<class B5, int, ONEAPI::bit_and<int>>(Q, ~0, 99, 32, 256);
-  testBoth<class B6, int, ONEAPI::minimum<int>>(
-      Q, (std::numeric_limits<int>::max)(), -99, 64, 256);
-  testBoth<class B7, int, ONEAPI::maximum<int>>(
-      Q, (std::numeric_limits<int>::min)(), 99, 128, 256);
-  testBoth<class B8, int, ONEAPI::plus<>>(Q, 0, 99, 256, 256);
+  tests<class B1, int>(Q, 0, 99, std::plus<>{}, 1, 32);
+  tests<class B2, int>(Q, 0, 99, std::bit_or<>{}, 8, 128);
+  tests<class B3, int>(Q, 0, 99, std::bit_xor<>{}, 16, 256);
+  tests<class B4, int>(Q, ~0, 99, std::bit_and<>{}, 32, 256);
+  tests<class B5, int>(Q, (std::numeric_limits<int>::max)(), -99,
+                       ext::oneapi::minimum<>{}, 64, 256);
+  tests<class B6, int>(Q, (std::numeric_limits<int>::min)(), 99,
+                       ext::oneapi::maximum<>{}, 128, 256);
+  tests<class B7, int>(Q, 0, 99, std::plus<>{}, 256, 256);
 
   // Check with various types.
-  testBoth<class C1, float, std::multiplies<float>>(Q, 1, 99, 8, 256);
-  testBoth<class C2, short, ONEAPI::minimum<>>(Q, 0x7fff, -99, 8, 256);
-  testBoth<class C3, unsigned char, ONEAPI::maximum<>>(Q, 0, 99, 8, 256);
+  tests<class C1, float>(Q, 1, 99, std::multiplies<>{}, 8, 24);
+  tests<class C2, short>(Q, 0x7fff, -99, ext::oneapi::minimum<>{}, 8, 256);
+  tests<class C3, unsigned char>(Q, 0, 99, ext::oneapi::maximum<>{}, 8, 256);
 
   // Check with CUSTOM type.
-  testBoth<class D1, CustomVec<long long>, CustomVecPlus<long long>>(
-      Q, CustomVec<long long>(0), CustomVec<long long>(99), 8, 256);
+  using CV = CustomVec<long long>;
+  tests<class D1, CV>(Q, CV(0), CV(99), CustomVecPlus<long long>{}, 8, 256);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_nd_s1_rw.cpp b/SYCL/Reduction/reduction_nd_s1_rw.cpp
index 2a40149555..654da943ec 100644
--- a/SYCL/Reduction/reduction_nd_s1_rw.cpp
+++ b/SYCL/Reduction/reduction_nd_s1_rw.cpp
@@ -7,104 +7,43 @@
 // with reductions initialized with 1-dimensional read_write accessor
 // accessing 1 element buffer.
 
-#include "reduction_utils.hpp"
-#include <CL/sycl.hpp>
-#include <cassert>
+#include "reduction_nd_range_scalar.hpp"
 
 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 T, bool B> class KName;
-
-template <typename Name, bool IsSYCL2020Mode, typename T, class BinaryOperation>
-void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  buffer<T, 1> InBuf(NWItems);
-  buffer<T, 1, CustomAllocator> OutBuf(1);
-
-  // Initialize.
-  BinaryOperation BOp;
-  T CorrectOut;
-  initInputData(InBuf, CorrectOut, Identity, BOp, NWItems);
-
-  // The final reduction sum after running parallel_for() must include
-  // the original value it was initialized with before the parallel_for().
-  CorrectOut = BOp(CorrectOut, Init);
-  (OutBuf.template get_access<access::mode::write>())[0] = Init;
-
-  // Compute.
-  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
-  if constexpr (IsSYCL2020Mode) {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      auto Redu = sycl::reduction(OutBuf, CGH, Identity, BOp);
-
-      CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-        Sum.combine(In[NDIt.get_global_linear_id()]);
-      });
-    });
-  } else {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      accessor<T, 1, access::mode::read_write, access::target::global_buffer>
-          Out(OutBuf, CGH);
-      auto Redu = ONEAPI::reduction(Out, Identity, BOp);
-
-      CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-        Sum.combine(In[NDIt.get_global_linear_id()]);
-      });
-    });
-  }
-
-  // Check correctness.
-  auto Out = OutBuf.template get_access<access::mode::read>();
-  T ComputedOut = *(Out.get_pointer());
-  if (ComputedOut != CorrectOut) {
-    std::cout << "NWItems = " << NWItems << ", WGSize = " << WGSize << "\n";
-    std::cout << "Computed value: " << ComputedOut
-              << ", Expected value: " << CorrectOut << "\n";
-    assert(0 && "Wrong value.");
-  }
-}
+int NumErrors = 0;
 
 template <typename Name, typename T, class BinaryOperation>
-void testBoth(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  test<KName<Name, false>, false, T, BinaryOperation>(Q, Identity, Init, WGSize,
-                                                      NWItems);
-  test<KName<Name, true>, true, T, BinaryOperation>(Q, Identity, Init, WGSize,
-                                                    NWItems);
+void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, size_t WGSize,
+           size_t NWItems) {
+  constexpr access::mode RW = access::mode::read_write;
+  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
+  NumErrors += testBoth<Name, RW>(Q, Identity, Init, BOp, NDRange);
 }
 
 int main() {
   queue Q;
+  printDeviceInfo(Q);
 
   // Check non power-of-two work-group sizes.
-  testBoth<class A1, int, ONEAPI::plus<int>>(Q, 0, 99, 1, 7);
-  testBoth<class A2, int, ONEAPI::plus<int>>(Q, 0, -99, 49, 49 * 5);
-
-  // Try some power-of-two work-group sizes.
-  testBoth<class B1, int, ONEAPI::plus<>>(Q, 0, 99, 2, 32);
-  testBoth<class B2, int, ONEAPI::plus<>>(Q, 0, 199, 32, 128);
-  testBoth<class B3, int, ONEAPI::plus<>>(Q, 0, 299, 128, 128);
-  testBoth<class B4, int, ONEAPI::plus<>>(Q, 0, 399, 256, 256);
+  tests<class A1, int>(Q, 0, 99, std::plus<int>{}, 1, 7);
+  tests<class A2, int>(Q, 0, -99, std::plus<int>{}, 49, 49 * 5);
 
   // Check with various operations and types.
-  testBoth<class C1, int, std::multiplies<int>>(Q, 1, 2, 8, 256);
-  testBoth<class C2, float, std::multiplies<float>>(Q, 1, 1.2, 8, 16);
-  testBoth<class C3, short, ONEAPI::bit_or<>>(Q, 0, 0x3400, 4, 32);
-  testBoth<class C4, int, ONEAPI::bit_xor<int>>(Q, 0, 0x12340000, 2, 16);
-  testBoth<class C5, char, ONEAPI::bit_and<>>(Q, ~0, ~0, 4, 16);
-  testBoth<class C6, int, ONEAPI::minimum<int>>(
-      Q, (std::numeric_limits<int>::max)(), 99, 8, 256);
-  testBoth<class C7, int, ONEAPI::maximum<float>>(
-      Q, (std::numeric_limits<int>::min)(), -99, 8, 256);
+  tests<class B1, int>(Q, 0, 99, std::plus<>{}, 1, 32);
+  tests<class B2, int>(Q, 0, 99, std::bit_or<>{}, 8, 128);
+  tests<class B3, int>(Q, 0, 99, std::bit_xor<>{}, 16, 256);
+  tests<class B4, int>(Q, ~0, 99, std::bit_and<>{}, 32, 256);
+  tests<class B5, int>(Q, (std::numeric_limits<int>::max)(), -99,
+                       ext::oneapi::minimum<>{}, 64, 256);
+  tests<class B6, int>(Q, (std::numeric_limits<int>::min)(), 99,
+                       ext::oneapi::maximum<>{}, 128, 256);
+  tests<class B7, int>(Q, 0, 99, std::plus<>{}, 256, 256);
 
   // Check with CUSTOM type.
-  testBoth<class D1, CustomVec<long long>, CustomVecPlus<long long>>(
-      Q, CustomVec<long long>(0), CustomVec<long long>(-199), 8, 256);
+  using CV = CustomVec<long long>;
+  tests<class C1, CV>(Q, CV(0), CV(-199), CustomVecPlus<long long>{}, 8, 256);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_placeholder.cpp b/SYCL/Reduction/reduction_placeholder.cpp
index e47b040497..f71bea0255 100644
--- a/SYCL/Reduction/reduction_placeholder.cpp
+++ b/SYCL/Reduction/reduction_placeholder.cpp
@@ -4,23 +4,24 @@
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
 // RUNx: %HOST_RUN_PLACEHOLDER %t.out
-// TODO: Enable the test for HOST when it supports ONEAPI::reduce() and
+// TODO: Enable the test for HOST when it supports ext::oneapi::reduce() and
 // barrier()
 
 // This test performs basic checks of parallel_for(nd_range, reduction, func)
 // with reductions initialized with a placeholder accessor.
 
 #include "reduction_utils.hpp"
-#include <CL/sycl.hpp>
-#include <cassert>
 
 using namespace cl::sycl;
 
-template <typename... Ts> class KNGroup;
+template <typename T, int N> class KName;
+
+template <typename Name, typename T, int Dim, class BinaryOperation,
+          access::mode Mode>
+int test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
+  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
+  printTestLabel<T, BinaryOperation>(false, NDRange);
 
-template <typename SpecializationKernelName, typename T, int Dim,
-          class BinaryOperation, access::mode Mode>
-void testOneCase(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
   // Initialize.
   T CorrectOut;
   BinaryOperation BOp;
@@ -39,57 +40,54 @@ void testOneCase(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
   Q.submit([&](handler &CGH) {
     auto In = InBuf.template get_access<access::mode::read>(CGH);
     CGH.require(Out);
-    auto Redu = ONEAPI::reduction(Out, Identity, BinaryOperation());
-    range<1> GlobalRange(NWItems);
-    range<1> LocalRange(WGSize);
-    nd_range<1> NDRange(GlobalRange, LocalRange);
-    CGH.parallel_for<SpecializationKernelName>(
-        NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
-          Sum.combine(In[NDIt.get_global_linear_id()]);
-        });
+    auto Redu = ext::oneapi::reduction(Out, Identity, BOp);
+    CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
+      Sum.combine(In[NDIt.get_global_linear_id()]);
+    });
   });
   Q.wait();
 
   // Check correctness.
   T ReduVar = (OutBuf.template get_access<access::mode::read>())[0];
-  if (ReduVar != CorrectOut) {
-    std::cout << "NWItems = " << NWItems << ", WGSize = " << WGSize << "\n";
-    std::cout << "Computed value: " << ReduVar
-              << ", Expected value: " << CorrectOut << "\n";
-    assert(0 && "Wrong value.");
-  }
+  return checkResults(Q, false /*SYCL2020*/, BOp, NDRange, ReduVar, CorrectOut);
 }
 
-template <typename KernelName, typename T, int Dim, class BinaryOperation>
-void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  testOneCase<KNGroup<KernelName, class RWCase>, T, Dim, BinaryOperation,
-              access::mode::read_write>(Q, Identity, Init, WGSize, NWItems);
-  testOneCase<KNGroup<KernelName, class DWCase>, T, Dim, BinaryOperation,
-              access::mode::discard_write>(Q, Identity, Init, WGSize, NWItems);
+int NumErrors = 0;
+
+template <typename Name, typename T, class BinaryOperation>
+void tests(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
+  constexpr access::mode DW = access::mode::discard_write;
+  constexpr access::mode RW = access::mode::read_write;
+
+  NumErrors += test<KName<Name, 0>, T, 0, BinaryOperation, DW>(
+      Q, Identity, Init, WGSize, NWItems);
+  NumErrors += test<KName<Name, 1>, T, 1, BinaryOperation, DW>(
+      Q, Identity, Init, WGSize, NWItems);
+
+  NumErrors += test<KName<Name, 2>, T, 0, BinaryOperation, RW>(
+      Q, Identity, Init, WGSize, NWItems);
+  NumErrors += test<KName<Name, 3>, T, 1, BinaryOperation, RW>(
+      Q, Identity, Init, WGSize, NWItems);
 }
 
 int main() {
   queue Q;
+  printDeviceInfo(Q);
+
   // fast atomics and fast reduce
-  test<class AtomicReduce1, int, 1, ONEAPI::plus<>>(Q, 0, 77, 49, 49 * 5);
-  test<class AtomicReduce2, int, 0, ONEAPI::plus<int>>(Q, 0, -77, 8, 8);
+  tests<class AtomicReduce, int, std::plus<>>(Q, 0, 77, 49, 49 * 5);
 
   // fast atomics
-  test<class Atomic1, int, 0, ONEAPI::bit_or<int>>(Q, 0, 233, 7, 7 * 3);
-  test<class Atomic2, int, 1, ONEAPI::bit_or<int>>(Q, 0, 177, 4, 128);
+  tests<class Atomic, int, std::bit_or<>>(Q, 0, 233, 7, 7 * 3);
 
   // fast reduce
-  test<class Reduce1, float, 1, ONEAPI::minimum<float>>(
+  tests<class Reduce, float, ext::oneapi::minimum<>>(
       Q, getMaximumFPValue<float>(), -5.0, 5, 5 * 7);
-  test<class Reduce2, float, 0, ONEAPI::maximum<float>>(
-      Q, getMinimumFPValue<float>(), -5.0, 4, 128);
 
   // generic algorithm
-  test<class Generic1, int, 0, std::multiplies<int>>(Q, 1, 2, 7, 7 * 5);
-  test<class Generic2, int, 1, std::multiplies<>>(Q, 1, 3, 8, 16);
-  test<class Generic3, CustomVec<short>, 0, CustomVecPlus<short>>(
+  tests<class Generic, CustomVec<short>, CustomVecPlus<short>>(
       Q, CustomVec<short>(0), CustomVec<short>(4), 8, 8 * 3);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_queue_parallel_for.cpp b/SYCL/Reduction/reduction_queue_parallel_for.cpp
index f143378fb5..ec3bfdd723 100644
--- a/SYCL/Reduction/reduction_queue_parallel_for.cpp
+++ b/SYCL/Reduction/reduction_queue_parallel_for.cpp
@@ -4,47 +4,38 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 
 // RUNx: %HOST_RUN_PLACEHOLDER %t.out
-// TODO: Enable the test for HOST when it supports ONEAPI::reduce() and
+// TODO: Enable the test for HOST when it supports ext::oneapi::reduce() and
 // barrier()
 
 // This test only checks that the method queue::parallel_for() accepting
 // reduction, can be properly translated into queue::submit + parallel_for().
 
-#include <CL/sycl.hpp>
+#include "reduction_utils.hpp"
+
 using namespace sycl;
 
 template <typename T, bool B> class KName;
 
-template <typename Name, bool IsSYCL2020Mode> int test(queue &Q) {
-  const size_t NElems = 1024;
-  const size_t WGSize = 256;
+template <typename Name, bool IsSYCL2020, typename T, typename BinaryOperation>
+int test(queue &Q, T Identity, size_t WGSize, size_t NElems) {
+  nd_range<1> NDRange(range<1>{NElems}, range<1>{WGSize});
+  printTestLabel<T, BinaryOperation>(IsSYCL2020, NDRange);
 
-  int *Data = malloc_shared<int>(NElems, Q);
+  T *Data = malloc_shared<T>(NElems, Q);
   for (int I = 0; I < NElems; I++)
     Data[I] = I;
 
-  int *Sum = malloc_shared<int>(1, Q);
-  *Sum = 0;
-
-  if constexpr (IsSYCL2020Mode) {
-    Q.parallel_for<Name>(
-         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<Name>(
-         nd_range<1>{NElems, WGSize}, ONEAPI::reduction(Sum, ONEAPI::plus<>()),
-         [=](nd_item<1> It, auto &Sum) { Sum += Data[It.get_global_id(0)]; })
-        .wait();
-  }
-
-  int ExpectedSum = (NElems - 1) * NElems / 2;
-  int Error = 0;
-  if (*Sum != ExpectedSum) {
-    std::cerr << "Error: Expected = " << ExpectedSum << ", Computed = " << *Sum
-              << std::endl;
-    Error = 1;
-  }
+  T *Sum = malloc_shared<T>(1, Q);
+  *Sum = Identity;
+
+  BinaryOperation BOp;
+  auto Redu = createReduction<IsSYCL2020, access::mode::read_write>(Sum, BOp);
+  Q.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> It, auto &Sum) {
+     Sum += Data[It.get_global_id(0)];
+   }).wait();
+
+  T ExpectedSum = (NElems - 1) * NElems / 2;
+  int Error = checkResults(Q, IsSYCL2020, BOp, NDRange, *Sum, ExpectedSum);
 
   free(Data, Q);
   free(Sum, Q);
@@ -53,7 +44,11 @@ template <typename Name, bool IsSYCL2020Mode> int test(queue &Q) {
 
 int main() {
   queue Q;
-  int Error = test<KName<class A, true>, true>(Q);
-  Error += test<KName<class A, false>, false>(Q);
-  return Error;
+  printDeviceInfo(Q);
+
+  int NumErrors = test<class A1, true, int, std::plus<>>(Q, 0, 16, 32);
+  NumErrors += test<class A2, false, int, std::plus<>>(Q, 0, 7, 14);
+
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_range_1d_s0_dw.cpp b/SYCL/Reduction/reduction_range_1d_s0_dw.cpp
index a000c9cfe7..e84d764ced 100644
--- a/SYCL/Reduction/reduction_range_1d_s0_dw.cpp
+++ b/SYCL/Reduction/reduction_range_1d_s0_dw.cpp
@@ -14,10 +14,13 @@
 
 using namespace cl::sycl;
 
+int NumErrors = 0;
+
 template <typename Name, typename T, class BinaryOperation>
 void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, size_t NWItems) {
   constexpr access::mode DW = access::mode::discard_write;
-  test<Name, false, DW, T, 0>(Q, Identity, Init, BOp, range<1>{NWItems});
+  NumErrors +=
+      test<Name, false, DW, 0>(Q, Identity, Init, BOp, range<1>{NWItems});
 }
 
 int main() {
@@ -38,19 +41,17 @@ int main() {
   tests<class B1, int>(Q, ~0, ~0, std::bit_and<>{}, 8);
   tests<class B2, int>(Q, 0, 0x12340000, std::bit_xor<>{}, 16);
   tests<class B3, int>(Q, 0, 0x3400, std::bit_or<>{}, MaxWGSize * 4);
-  tests<class B4, int>(Q, 1, 2, std::multiplies<>{}, 256);
-  tests<class B5, int>(Q, 1, 3, std::multiplies<>{}, MaxWGSize + 1);
-  tests<class B6, int>(Q, (std::numeric_limits<int>::max)(), -99,
+  tests<class B4, uint64_t>(Q, 1, 2, std::multiplies<>{}, 31);
+  tests<class B5, int>(Q, (std::numeric_limits<int>::max)(), -99,
                        ext::oneapi::minimum<>{}, MaxWGSize * 2);
-  tests<class B7, int>(Q, (std::numeric_limits<int>::min)(), 99,
+  tests<class B6, int>(Q, (std::numeric_limits<int>::min)(), 99,
                        ext::oneapi::maximum<>{}, 8);
-  tests<class B8, uint64_t>(Q, 1, 99, std::multiplies<>{}, 37);
 
   // Check with CUSTOM type.
   using CV = CustomVec<long long>;
   tests<class C1>(Q, CV(0), CV(99), CustomVecPlus<long long>{}, 64);
   tests<class C2>(Q, CV(0), CV(99), CustomVecPlus<long long>{}, MaxWGSize * 3);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_range_1d_s0_rw.cpp b/SYCL/Reduction/reduction_range_1d_s0_rw.cpp
index 8e23e9e26a..bc1e2befed 100644
--- a/SYCL/Reduction/reduction_range_1d_s0_rw.cpp
+++ b/SYCL/Reduction/reduction_range_1d_s0_rw.cpp
@@ -14,10 +14,13 @@
 
 using namespace cl::sycl;
 
+int NumErrors = 0;
+
 template <typename Name, typename T, class BinaryOperation>
 void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, size_t NWItems) {
   constexpr access::mode RW = access::mode::read_write;
-  test<Name, false, RW, T, 0>(Q, Identity, Init, BOp, range<1>{NWItems});
+  NumErrors +=
+      test<Name, false, RW, 0>(Q, Identity, Init, BOp, range<1>{NWItems});
 }
 
 int main() {
@@ -38,13 +41,12 @@ int main() {
   tests<class B1, int>(Q, ~0, ~0, std::bit_and<>{}, 8);
   tests<class B2, int>(Q, 0, 0x12340000, std::bit_xor<>{}, 16);
   tests<class B3, int>(Q, 0, 0x3400, std::bit_or<>{}, MaxWGSize * 4);
-  tests<class B4, int>(Q, 1, 2, std::multiplies<>{}, 256);
-  tests<class B5, int>(Q, 1, 3, std::multiplies<>{}, MaxWGSize + 1);
-  tests<class B6, int>(Q, (std::numeric_limits<int>::max)(), -99,
+  tests<class B4, uint64_t>(Q, 1, 2, std::multiplies<>{}, 17);
+  tests<class B5, int>(Q, (std::numeric_limits<int>::max)(), -99,
                        ext::oneapi::minimum<>{}, MaxWGSize * 2);
-  tests<class B7, int>(Q, (std::numeric_limits<int>::min)(), 99,
+  tests<class B6, int>(Q, (std::numeric_limits<int>::min)(), 99,
                        ext::oneapi::maximum<>{}, 8);
-  tests<class B8, float>(Q, 1, 99, std::multiplies<>{}, MaxWGSize);
+  tests<class B7, float>(Q, 1, 99, std::multiplies<>{}, 16);
 
   // Check with CUSTOM type.
   tests<class C1>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
@@ -52,6 +54,6 @@ int main() {
   tests<class C2>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
                   CustomVecPlus<long long>{}, MaxWGSize * 3);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_range_1d_s1_dw.cpp b/SYCL/Reduction/reduction_range_1d_s1_dw.cpp
index e2924a0050..ff5c01898e 100644
--- a/SYCL/Reduction/reduction_range_1d_s1_dw.cpp
+++ b/SYCL/Reduction/reduction_range_1d_s1_dw.cpp
@@ -14,10 +14,12 @@
 
 using namespace cl::sycl;
 
+int NumErrors = 0;
+
 template <typename Name, typename T, class BinaryOperation>
 void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, size_t NWItems) {
   constexpr access::mode DW = access::mode::discard_write;
-  testBoth<Name, DW, T>(Q, Identity, Init, BOp, range<1>{NWItems});
+  NumErrors += testBoth<Name, DW>(Q, Identity, Init, BOp, range<1>{NWItems});
 }
 
 int main() {
@@ -27,24 +29,22 @@ int main() {
       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<>{}, 1);
-  tests<class A2, int>(Q, 0, 99, std::plus<>{}, 2);
-  tests<class A3, int64_t>(Q, 0, 99, std::plus<>{}, 7);
-  tests<class A4, int64_t>(Q, 0, 99, std::plus<>{}, 64);
-  tests<class A5, int>(Q, 0, 99, std::plus<>{}, MaxWGSize * 2);
-  tests<class A6, int>(Q, 0, 99, std::plus<>{}, MaxWGSize * 2 + 5);
+  tests<class A1, int>(Q, 0, 99, std::plus<int>{}, 1);
+  tests<class A2, int64_t>(Q, 0, 99, std::plus<>{}, 7);
+  tests<class A3, int64_t>(Q, 0, 99, std::plus<>{}, 64);
+  tests<class A4, int>(Q, 0, 99, std::plus<>{}, MaxWGSize * 2);
+  tests<class A5, int>(Q, 0, 99, std::plus<>{}, MaxWGSize * 2 + 5);
 
   // Try various types & ranges.
   tests<class B1, int>(Q, ~0, 99, std::bit_and<>{}, 7);
   tests<class B2, int>(Q, 0, 0xff99, std::bit_xor<>{}, MaxWGSize);
   tests<class B3, int>(Q, 0, 0xff99, std::bit_or<>{}, 3);
-  tests<class B4, int>(Q, 1, 3, std::multiplies<>{}, 32);
-  tests<class B5, int>(Q, 1, 3, std::multiplies<>{}, MaxWGSize * 4);
-  tests<class B6, int>(Q, (std::numeric_limits<int>::max)(), -99,
+  tests<class B4, short>(Q, 1, 2, std::multiplies<>{}, 7);
+  tests<class B5, int>(Q, (std::numeric_limits<int>::max)(), -99,
                        ext::oneapi::minimum<>{}, MaxWGSize * 2);
-  tests<class B7, int>(Q, (std::numeric_limits<int>::min)(), 99,
+  tests<class B6, int>(Q, (std::numeric_limits<int>::min)(), 99,
                        ext::oneapi::maximum<>{}, 8);
-  tests<class B8, float>(Q, 1, 99, std::multiplies<>{}, MaxWGSize);
+  tests<class B7, float>(Q, 1, 99, std::multiplies<>{}, 10);
 
   // Check with CUSTOM type.
   tests<class C1>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
@@ -52,6 +52,6 @@ int main() {
   tests<class C2>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
                   CustomVecPlus<long long>{}, MaxWGSize * 3);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_range_1d_s1_rw.cpp b/SYCL/Reduction/reduction_range_1d_s1_rw.cpp
index 01d55587b7..bc10d647ce 100644
--- a/SYCL/Reduction/reduction_range_1d_s1_rw.cpp
+++ b/SYCL/Reduction/reduction_range_1d_s1_rw.cpp
@@ -18,10 +18,12 @@
 
 using namespace cl::sycl;
 
+int NumErrors = 0;
+
 template <typename Name, typename T, class BinaryOperation>
 void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, size_t NWItems) {
   constexpr access::mode RW = access::mode::read_write;
-  testBoth<Name, RW, T>(Q, Identity, Init, BOp, range<1>{NWItems});
+  NumErrors += testBoth<Name, RW>(Q, Identity, Init, BOp, range<1>{NWItems});
 }
 
 int main() {
@@ -32,23 +34,21 @@ int main() {
 
   // Fast-reduce and Fast-atomics. Try various range types/sizes.
   tests<class A1, int>(Q, 0, 99, std::plus<int>{}, 1);
-  tests<class A2, int>(Q, 0, 99, std::plus<int>{}, 2);
-  tests<class A3, int64_t>(Q, 0, 99, std::plus<>{}, 7);
-  tests<class A4, int64_t>(Q, 0, 99, std::plus<>{}, 64);
-  tests<class A5, int>(Q, 0, 99, std::plus<>{}, MaxWGSize * 2);
-  tests<class A6, int>(Q, 0, 99, std::plus<>{}, MaxWGSize * 2 + 5);
+  tests<class A2, int64_t>(Q, 0, 99, std::plus<>{}, 7);
+  tests<class A3, int64_t>(Q, 0, 99, std::plus<>{}, 64);
+  tests<class A4, int>(Q, 0, 99, std::plus<>{}, MaxWGSize * 2);
+  tests<class A5, int>(Q, 0, 99, std::plus<>{}, MaxWGSize * 2 + 5);
 
   // Try various types & ranges.
   tests<class B1, int>(Q, ~0, ~0, std::bit_and<>{}, 8);
   tests<class B2, int>(Q, 0, 0x12340000, std::bit_xor<>{}, 16);
   tests<class B3, int>(Q, 0, 0x3400, std::bit_or<>{}, MaxWGSize * 4);
-  tests<class B4, int>(Q, 1, 2, std::multiplies<>{}, 256);
-  tests<class B5, int>(Q, 1, 3, std::multiplies<>{}, MaxWGSize * 4);
+  tests<class B4, uint64_t>(Q, 1, 2, std::multiplies<>{}, 16);
+  tests<class B5, float>(Q, 1, 3, std::multiplies<>{}, 11);
   tests<class B6, int>(Q, (std::numeric_limits<int>::max)(), -99,
                        ext::oneapi::minimum<>{}, MaxWGSize * 2);
   tests<class B7, int>(Q, (std::numeric_limits<int>::min)(), 99,
                        ext::oneapi::maximum<>{}, 8);
-  tests<class B8, float>(Q, 1, 99, std::multiplies<>{}, MaxWGSize);
 
   // Check with CUSTOM type.
   tests<class C1>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
@@ -56,6 +56,6 @@ int main() {
   tests<class C2>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
                   CustomVecPlus<long long>{}, MaxWGSize * 3);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_range_2d_s1_dw.cpp b/SYCL/Reduction/reduction_range_2d_s1_dw.cpp
index 0846b0ac0e..79476a1f52 100644
--- a/SYCL/Reduction/reduction_range_2d_s1_dw.cpp
+++ b/SYCL/Reduction/reduction_range_2d_s1_dw.cpp
@@ -14,10 +14,12 @@
 
 using namespace cl::sycl;
 
+int NumErrors = 0;
+
 template <typename Name, typename T, class BinaryOperation>
 void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, range<2> Range) {
   constexpr access::mode DW = access::mode::discard_write;
-  testBoth<Name, DW, T>(Q, Identity, Init, BOp, Range);
+  NumErrors += testBoth<Name, DW>(Q, Identity, Init, BOp, Range);
 }
 
 int main() {
@@ -38,9 +40,8 @@ int main() {
   tests<class B1, int>(Q, 0, 0x2021ff99, std::bit_xor<>{}, range<2>{3, 3});
   tests<class B2, int>(Q, ~0, 99, std::bit_and<>{}, range<2>{4, 3});
   tests<class B3, int>(Q, 0, 99, std::bit_or<>{}, range<2>{2, 2});
-  tests<class B4, uint64_t>(Q, 1, 3, std::multiplies<>{}, range<2>{16, 3});
-  tests<class B5, uint64_t>(Q, 1, 3, std::multiplies<>{},
-                            range<2>{3, MaxWGSize});
+  tests<class B4, uint64_t>(Q, 1, 3, std::multiplies<>{}, range<2>{8, 3});
+  tests<class B5, uint64_t>(Q, 1, 3, std::multiplies<>{}, range<2>{3, 7});
   tests<class B6, int>(Q, (std::numeric_limits<int>::max)(), -99,
                        ext::oneapi::minimum<>{}, range<2>{8, 3});
   tests<class B7, int>(Q, (std::numeric_limits<int>::min)(), 99,
@@ -50,6 +51,6 @@ int main() {
   tests<class C1>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
                   CustomVecPlus<long long>{}, range<2>{33, MaxWGSize});
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_range_2d_s1_rw.cpp b/SYCL/Reduction/reduction_range_2d_s1_rw.cpp
index 56a5a9e7af..4dea17a2b7 100644
--- a/SYCL/Reduction/reduction_range_2d_s1_rw.cpp
+++ b/SYCL/Reduction/reduction_range_2d_s1_rw.cpp
@@ -14,10 +14,12 @@
 
 using namespace cl::sycl;
 
+int NumErrors = 0;
+
 template <typename Name, typename T, class BinaryOperation>
 void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, range<2> Range) {
   constexpr access::mode RW = access::mode::read_write;
-  testBoth<Name, RW, T>(Q, Identity, Init, BOp, Range);
+  NumErrors += testBoth<Name, RW>(Q, Identity, Init, BOp, Range);
 }
 
 int main() {
@@ -38,9 +40,8 @@ int main() {
   tests<class B1, int>(Q, 0, 0x2021ff99, std::bit_xor<>{}, range<2>{3, 3});
   tests<class B2, int>(Q, ~0, 99, std::bit_and<>{}, range<2>{4, 3});
   tests<class B3, int>(Q, 0, 99, std::bit_or<>{}, range<2>{2, 2});
-  tests<class B4, uint64_t>(Q, 1, 3, std::multiplies<>{}, range<2>{16, 3});
-  tests<class B5, uint64_t>(Q, 1, 3, std::multiplies<>{},
-                            range<2>{3, MaxWGSize});
+  tests<class B4, uint64_t>(Q, 1, 3, std::multiplies<>{}, range<2>{8, 3});
+  tests<class B5, uint64_t>(Q, 1, 3, std::multiplies<>{}, range<2>{3, 8});
   tests<class B6, int>(Q, (std::numeric_limits<int>::max)(), -99,
                        ext::oneapi::minimum<>{}, range<2>{8, 3});
   tests<class B7, int>(Q, (std::numeric_limits<int>::min)(), 99,
@@ -50,6 +51,6 @@ int main() {
   tests<class C1>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
                   CustomVecPlus<long long>{}, range<2>{33, MaxWGSize});
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_range_3d_s1_dw.cpp b/SYCL/Reduction/reduction_range_3d_s1_dw.cpp
index d0b11a98f5..26fb3b1c50 100644
--- a/SYCL/Reduction/reduction_range_3d_s1_dw.cpp
+++ b/SYCL/Reduction/reduction_range_3d_s1_dw.cpp
@@ -14,10 +14,12 @@
 
 using namespace cl::sycl;
 
+int NumErrors = 0;
+
 template <typename Name, typename T, class BinaryOperation>
 void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, range<3> Range) {
   constexpr access::mode DW = access::mode::discard_write;
-  testBoth<Name, DW, T>(Q, Identity, Init, BOp, Range);
+  NumErrors += testBoth<Name, DW>(Q, Identity, Init, BOp, Range);
 }
 
 int main() {
@@ -50,19 +52,17 @@ int main() {
                        range<3>{MaxWGSize * 3, 4, 3});
   tests<class B3, int>(Q, 0, 99, std::bit_or<>{},
                        range<3>{2, 2, MaxWGSize * 3});
-  tests<class B4, uint64_t>(Q, 1, 3, std::multiplies<>{}, range<3>{16, 3, 5});
-  tests<class B5, uint64_t>(Q, 1, 3, std::multiplies<>{},
-                            range<3>{2, 3, MaxWGSize});
+  tests<class B4, uint64_t>(Q, 1, 3, std::multiplies<>{}, range<3>{2, 3, 5});
+  tests<class B5, uint64_t>(Q, 1, 2, std::multiplies<>{}, range<3>{2, 3, 8});
   tests<class B6, int>(Q, (std::numeric_limits<int>::max)(), -99,
                        ext::oneapi::minimum<>{}, range<3>{MaxWGSize, 8, 3});
   tests<class B7, int>(Q, (std::numeric_limits<int>::min)(), 99,
                        ext::oneapi::maximum<>{}, range<3>{3, MaxWGSize, 3});
-  tests<class B8, float>(Q, 1, 99, std::multiplies<>{},
-                         range<3>{3, 3, MaxWGSize});
+  tests<class B8, float>(Q, 1, 99, std::multiplies<>{}, range<3>{3, 3, 5});
 
   tests<class C1>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
                   CustomVecPlus<long long>{}, range<3>{2, 33, MaxWGSize});
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_range_3d_s1_rw.cpp b/SYCL/Reduction/reduction_range_3d_s1_rw.cpp
index 963e9b4769..b5b8ddc0d3 100644
--- a/SYCL/Reduction/reduction_range_3d_s1_rw.cpp
+++ b/SYCL/Reduction/reduction_range_3d_s1_rw.cpp
@@ -14,10 +14,12 @@
 
 using namespace cl::sycl;
 
+int NumErrors = 0;
+
 template <typename Name, typename T, class BinaryOperation>
 void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, range<3> Range) {
   constexpr access::mode RW = access::mode::read_write;
-  testBoth<Name, RW, T>(Q, Identity, Init, BOp, Range);
+  NumErrors += testBoth<Name, RW>(Q, Identity, Init, BOp, Range);
 }
 
 int main() {
@@ -50,19 +52,17 @@ int main() {
                        range<3>{MaxWGSize * 3, 4, 3});
   tests<class B3, int>(Q, 0, 99, std::bit_or<>{},
                        range<3>{2, 2, MaxWGSize * 3});
-  tests<class B4, uint64_t>(Q, 1, 3, std::multiplies<>{}, range<3>{16, 3, 5});
-  tests<class B5, uint64_t>(Q, 1, 3, std::multiplies<>{},
-                            range<3>{2, 3, MaxWGSize});
+  tests<class B4, uint64_t>(Q, 1, 3, std::multiplies<>{}, range<3>{2, 3, 5});
+  tests<class B5, uint64_t>(Q, 1, 3, std::multiplies<>{}, range<3>{5, 2, 3});
   tests<class B6, int>(Q, (std::numeric_limits<int>::max)(), -99,
                        ext::oneapi::minimum<>{}, range<3>{MaxWGSize, 8, 3});
   tests<class B7, int>(Q, (std::numeric_limits<int>::min)(), 99,
                        ext::oneapi::maximum<>{}, range<3>{3, MaxWGSize, 3});
-  tests<class B8, float>(Q, 1, 99, std::multiplies<>{},
-                         range<3>{3, 3, MaxWGSize});
+  tests<class B8, float>(Q, 1, 99, std::multiplies<>{}, range<3>{3, 3, 4});
 
   tests<class C1>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
                   CustomVecPlus<long long>{}, range<3>{2, 33, MaxWGSize});
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_range_lambda.cpp b/SYCL/Reduction/reduction_range_lambda.cpp
index ef7a1d46be..b677055609 100644
--- a/SYCL/Reduction/reduction_range_lambda.cpp
+++ b/SYCL/Reduction/reduction_range_lambda.cpp
@@ -25,26 +25,34 @@ int main() {
 
   auto LambdaSum = [](auto x, auto y) { return (x + y); };
 
-  testBoth<class A1, RW, int>(Q, 0, 99, LambdaSum, range<1>{7});
-  testBoth<class A2, DW, int>(Q, 0, 99, LambdaSum, range<1>{7});
+  int NumErrors = 0;
 
-  testBoth<class A3, RW, int>(Q, 0, 99, LambdaSum, range<1>{MaxWGSize + 1});
-  testBoth<class A4, DW, int>(Q, 0, 99, LambdaSum, range<1>{MaxWGSize + 1});
+  NumErrors += testBoth<class A1, RW, int>(Q, 0, 99, LambdaSum, range<1>{7});
+  NumErrors += testBoth<class A2, DW, int>(Q, 0, 99, LambdaSum, range<1>{7});
 
-  testBoth<class B1, RW, int>(Q, 0, 99, LambdaSum, range<2>{3, 4});
-  testBoth<class B2, DW, int>(Q, 0, 99, LambdaSum, range<2>{3, 4});
+  NumErrors +=
+      testBoth<class A3, RW, int>(Q, 0, 99, LambdaSum, range<1>{MaxWGSize + 1});
+  NumErrors +=
+      testBoth<class A4, DW, int>(Q, 0, 99, LambdaSum, range<1>{MaxWGSize + 1});
 
-  testBoth<class B3, RW, int>(Q, 0, 99, LambdaSum, range<2>{3, MaxWGSize + 1});
-  testBoth<class B4, DW, int>(Q, 0, 99, LambdaSum, range<2>{3, MaxWGSize + 1});
+  NumErrors += testBoth<class B1, RW, int>(Q, 0, 99, LambdaSum, range<2>{3, 4});
+  NumErrors += testBoth<class B2, DW, int>(Q, 0, 99, LambdaSum, range<2>{3, 4});
 
-  testBoth<class C1, RW, int>(Q, 0, 99, LambdaSum, range<3>{2, 3, 4});
-  testBoth<class C2, RW, int>(Q, 0, 99, LambdaSum, range<3>{2, 3, 4});
+  NumErrors += testBoth<class B3, RW, int>(Q, 0, 99, LambdaSum,
+                                           range<2>{3, MaxWGSize + 1});
+  NumErrors += testBoth<class B4, DW, int>(Q, 0, 99, LambdaSum,
+                                           range<2>{3, MaxWGSize + 1});
 
-  testBoth<class C3, RW, int>(Q, 0, 99, LambdaSum,
-                              range<3>{2, 3, MaxWGSize + 1});
-  testBoth<class C4, DW, int>(Q, 0, 99, LambdaSum,
-                              range<3>{2, 3, MaxWGSize + 1});
+  NumErrors +=
+      testBoth<class C1, RW, int>(Q, 0, 99, LambdaSum, range<3>{2, 3, 4});
+  NumErrors +=
+      testBoth<class C2, RW, int>(Q, 0, 99, LambdaSum, range<3>{2, 3, 4});
 
-  std::cout << "Test passed\n";
-  return 0;
+  NumErrors += testBoth<class C3, RW, int>(Q, 0, 99, LambdaSum,
+                                           range<3>{2, 3, MaxWGSize + 1});
+  NumErrors += testBoth<class C4, DW, int>(Q, 0, 99, LambdaSum,
+                                           range<3>{2, 3, MaxWGSize + 1});
+
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_range_scalar.hpp b/SYCL/Reduction/reduction_range_scalar.hpp
index 7ee833f9d1..ef03f2ace1 100644
--- a/SYCL/Reduction/reduction_range_scalar.hpp
+++ b/SYCL/Reduction/reduction_range_scalar.hpp
@@ -3,38 +3,21 @@
 // accessing a scalar holding the reduction result.
 
 #include "reduction_utils.hpp"
-#include <CL/sycl.hpp>
-#include <cassert>
 
 using namespace cl::sycl;
 
 template <typename T, bool B> class KName;
 
-template <int Dims>
-std::ostream &operator<<(std::ostream &OS, const range<Dims> &Range) {
-  OS << "{" << Range[0];
-  if constexpr (Dims > 1)
-    OS << ", " << Range[1];
-  if constexpr (Dims > 2)
-    OS << ", " << Range[2];
-  OS << "}";
-  return OS;
-}
-
-template <typename Name, bool IsSYCL2020Mode, access::mode Mode, typename T,
-          int AccDim = 1, class BinaryOperation, int Dims>
-void test(queue &Q, T Identity, T Init, BinaryOperation BOp,
-          range<Dims> Range) {
-  std::string StdMode = IsSYCL2020Mode ? "SYCL2020" : "ONEAPI  ";
-  std::cout << "Running the test case: " << StdMode
-            << " {T=" << typeid(T).name()
-            << ", BOp=" << typeid(BinaryOperation).name() << ", Range=" << Range
-            << std::endl;
+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)
-    return;
+    return 0;
 
   buffer<T, Dims> InBuf(Range);
   buffer<T, 1> OutBuf(1);
@@ -51,42 +34,24 @@ void test(queue &Q, T Identity, T Init, BinaryOperation BOp,
   (OutBuf.template get_access<access::mode::write>())[0] = Init;
 
   // Compute.
-  if constexpr (IsSYCL2020Mode) {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      property_list PropList = getPropertyList<Mode>();
-      auto Redu = sycl::reduction(OutBuf, CGH, Identity, BOp, PropList);
-
-      CGH.parallel_for<Name>(
-          Range, Redu, [=](id<Dims> Id, auto &Sum) { Sum.combine(In[Id]); });
-    });
-  } else {
-    Q.submit([&](handler &CGH) {
-      auto In = InBuf.template get_access<access::mode::read>(CGH);
-      accessor<T, AccDim, Mode, access::target::global_buffer> Out(OutBuf, CGH);
-      auto Redu = ext::oneapi::reduction(Out, Identity, BOp);
-
-      CGH.parallel_for<Name>(
-          Range, Redu, [=](id<Dims> Id, auto &Sum) { Sum.combine(In[Id]); });
-    });
-  }
+  Q.submit([&](handler &CGH) {
+    auto In = InBuf.template get_access<access::mode::read>(CGH);
+    auto Redu =
+        createReduction<IsSYCL2020, Mode, AccDim>(OutBuf, CGH, Identity, BOp);
+    CGH.parallel_for<Name>(
+        Range, Redu, [=](id<Dims> Id, auto &Sum) { Sum.combine(In[Id]); });
+  });
 
   // Check correctness.
   auto Out = OutBuf.template get_access<access::mode::read>();
   T ComputedOut = *(Out.get_pointer());
-  if (ComputedOut != CorrectOut) {
-    printDeviceInfo(Q, true);
-    std::cerr << "Error: Range = " << Range << ", "
-              << "Computed value: " << ComputedOut
-              << ", Expected value: " << CorrectOut << "\n";
-    assert(0 && "Wrong value.");
-  }
+  return checkResults(Q, IsSYCL2020, BOp, Range, ComputedOut, CorrectOut);
 }
 
 template <typename Name, access::mode Mode, typename T, class BinaryOperation,
           int Dims>
-void testBoth(queue &Q, T Identity, T Init, BinaryOperation BOp,
-              range<Dims> Range) {
-  test<KName<Name, false>, false, Mode>(Q, Identity, Init, BOp, Range);
-  test<KName<Name, true>, true, Mode>(Q, Identity, Init, BOp, Range);
+int testBoth(queue &Q, T Identity, T Init, BinaryOperation BOp,
+             const range<Dims> &Range) {
+  return test<KName<Name, false>, false, Mode>(Q, Identity, Init, BOp, Range) +
+         test<KName<Name, true>, true, Mode>(Q, Identity, Init, BOp, Range);
 }
diff --git a/SYCL/Reduction/reduction_reducer_op_eq.cpp b/SYCL/Reduction/reduction_reducer_op_eq.cpp
index 07c4bbce28..e6990c875e 100644
--- a/SYCL/Reduction/reduction_reducer_op_eq.cpp
+++ b/SYCL/Reduction/reduction_reducer_op_eq.cpp
@@ -6,9 +6,7 @@
 // This test checks that operators ++, +=, *=, |=, &=, ^= are supported
 // whent the corresponding std::plus<>, std::multiplies, etc are defined.
 
-#include <CL/sycl.hpp>
-#include <cmath>
-#include <iostream>
+#include "reduction_utils.hpp"
 
 using namespace sycl;
 
@@ -78,19 +76,13 @@ template <> struct bit_and<XY> {
 };
 } // namespace std
 
-template <bool IsSYCL2020Mode, typename T, typename BinaryOperation>
-auto createReduction(T *USMPtr, T Identity, BinaryOperation BOp) {
-  if constexpr (IsSYCL2020Mode)
-    return sycl::reduction(USMPtr, Identity, BOp);
-  else
-    return ONEAPI::reduction(USMPtr, Identity, BOp);
-}
-
 template <typename T, bool IsSYCL2020Mode, typename BinaryOperation,
           OperationEqual OpEq, bool IsFP>
 int test(queue &Q, T Identity) {
   constexpr size_t N = 16;
   constexpr size_t L = 4;
+  nd_range<1> NDR{N, L};
+  printTestLabel<T, BinaryOperation>(IsSYCL2020Mode, NDR);
 
   T *Data = malloc_host<T>(N, Q);
   T *Res = malloc_host<T>(1, Q);
@@ -107,7 +99,6 @@ int test(queue &Q, T Identity) {
 
   *Res = Identity;
   auto Red = createReduction<IsSYCL2020Mode>(Res, Identity, BOp);
-  nd_range<1> NDR{N, L};
   if constexpr (OpEq == PlusPlus) {
     auto Lambda = [=](nd_item<1> ID, auto &Sum) { ++Sum; };
     Q.submit([&](handler &H) { H.parallel_for(NDR, Red, Lambda); }).wait();
@@ -189,11 +180,12 @@ template <typename T, bool TestPlusPlus> int testINTPack(queue &Q) {
 
 int main() {
   queue Q;
-  int Error = 0;
-  Error += testFPPack<float2>(Q);
-  Error += testINTPack<int2, true>(Q);
-  Error += testINTPack<XY, false>(Q);
-
-  std::cout << (Error ? "Failed\n" : "Passed.\n");
-  return Error;
+  printDeviceInfo(Q);
+  int NumErrors = 0;
+  NumErrors += testFPPack<float2>(Q);
+  NumErrors += testINTPack<int2, true>(Q);
+  NumErrors += testINTPack<XY, false>(Q);
+
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_usm.cpp b/SYCL/Reduction/reduction_usm.cpp
index 941c113881..e3c22b17d2 100644
--- a/SYCL/Reduction/reduction_usm.cpp
+++ b/SYCL/Reduction/reduction_usm.cpp
@@ -4,46 +4,33 @@
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
 // This test performs basic checks of parallel_for(nd_range, reduction, func)
-// with reductions initialized with USM var. It tests both ONEAPI::reduction
-// and SYCL-2020 reduction (sycl::reduction) assuming only read-write access,
-// i.e. without using SYCL-2020 property::reduction::initialize_to_identity.
+// with reductions initialized with USM var. It tests both
+// ext::oneapi::reduction and SYCL-2020 reduction (sycl::reduction) assuming
+// only read-write access, i.e. without using SYCL-2020
+// property::reduction::initialize_to_identity.
 
 #include "reduction_utils.hpp"
-#include <CL/sycl.hpp>
-#include <cassert>
 
 using namespace cl::sycl;
 
-template <typename T1, typename T2> class KernelNameGroup;
-
-template <bool IsSYCL2020, typename T, typename BinaryOperation>
-auto createReduction(T *USMPtr, T Identity, BinaryOperation BOp) {
-  if constexpr (IsSYCL2020)
-    return sycl::reduction(USMPtr, Identity, BOp);
-  else
-    return ONEAPI::reduction(USMPtr, Identity, BOp);
-}
+template <typename T1, typename T2> class KName;
 
 template <typename Name, bool IsSYCL2020, typename T, class BinaryOperation>
-void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems,
-          usm::alloc AllocType) {
+int test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems,
+         usm::alloc AllocType) {
+  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
+  printTestLabel<T, BinaryOperation>(IsSYCL2020, NDRange);
+
   auto Dev = Q.get_device();
-  if (AllocType == usm::alloc::shared &&
-      !Dev.get_info<info::device::usm_shared_allocations>())
-    return;
-  if (AllocType == usm::alloc::host &&
-      !Dev.get_info<info::device::usm_host_allocations>())
-    return;
-  if (AllocType == usm::alloc::device &&
-      !Dev.get_info<info::device::usm_device_allocations>())
-    return;
+  if (!Dev.has(getUSMAspect(AllocType)))
+    return 0;
 
   T *ReduVarPtr = (T *)malloc(sizeof(T), Dev, Q.get_context(), AllocType);
   if (ReduVarPtr == nullptr)
-    return;
+    return 0;
   if (AllocType == usm::alloc::device) {
     Q.submit([&](handler &CGH) {
-       CGH.single_task<KernelNameGroup<Name, class InitKernel>>(
+       CGH.single_task<KName<Name, class InitKernel>>(
            [=]() { *ReduVarPtr = Init; });
      }).wait();
   } else {
@@ -62,9 +49,7 @@ void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems,
   Q.submit([&](handler &CGH) {
      auto In = InBuf.template get_access<access::mode::read>(CGH);
      auto Redu = createReduction<IsSYCL2020>(ReduVarPtr, Identity, BOp);
-     nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
-
-     CGH.parallel_for<KernelNameGroup<Name, class Test>>(
+     CGH.parallel_for<KName<Name, class Test>>(
          NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
            Sum.combine(In[NDIt.get_global_linear_id()]);
          });
@@ -76,63 +61,65 @@ void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems,
     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<KernelNameGroup<Name, class Check>>(
+       CGH.single_task<KName<Name, class Check>>(
            [=]() { OutAcc[0] = *ReduVarPtr; });
      }).wait();
     ComputedOut = (Buf.template get_access<access::mode::read>())[0];
   } else {
     ComputedOut = *ReduVarPtr;
   }
-  if (ComputedOut != CorrectOut) {
-    std::cerr << "NWItems = " << NWItems << ", WGSize = " << WGSize << "\n";
-    std::cerr << "Computed value: " << ComputedOut
-              << ", Expected value: " << CorrectOut
-              << ", AllocMode: " << static_cast<int>(AllocType) << "\n";
-    assert(0 && "Wrong value.");
-  }
 
+  std::string AllocStr =
+      "AllocMode=" + std::to_string(static_cast<int>(AllocType));
+  int Error = checkResults(Q, IsSYCL2020, BOp, NDRange, ComputedOut, CorrectOut,
+                           AllocStr);
   free(ReduVarPtr, Q.get_context());
+  return Error;
 }
 
+int NumErrors = 0;
+
 template <typename Name, typename T, class BinaryOperation>
 void testUSM(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
   // Test SYCL-2020 reductions
-  test<KernelNameGroup<Name, class Shared2020>, true, T, BinaryOperation>(
+  NumErrors += test<KName<Name, class Shared2020>, true, T, BinaryOperation>(
       Q, Identity, Init, WGSize, NWItems, usm::alloc::shared);
-  test<KernelNameGroup<Name, class Host2020>, true, T, BinaryOperation>(
+  NumErrors += test<KName<Name, class Host2020>, true, T, BinaryOperation>(
       Q, Identity, Init, WGSize, NWItems, usm::alloc::host);
-  test<KernelNameGroup<Name, class Device2020>, true, T, BinaryOperation>(
+  NumErrors += test<KName<Name, class Device2020>, true, T, BinaryOperation>(
       Q, Identity, Init, WGSize, NWItems, usm::alloc::device);
 
-  // Test ONEAPI reductions
-  test<KernelNameGroup<Name, class Shared>, false, T, BinaryOperation>(
+  // Test ext::oneapi reductions
+  NumErrors += test<KName<Name, class Shared>, false, T, BinaryOperation>(
       Q, Identity, Init, WGSize, NWItems, usm::alloc::shared);
-  test<KernelNameGroup<Name, class Host>, false, T, BinaryOperation>(
+  NumErrors += test<KName<Name, class Host>, false, T, BinaryOperation>(
       Q, Identity, Init, WGSize, NWItems, usm::alloc::host);
-  test<KernelNameGroup<Name, class Device>, false, T, BinaryOperation>(
+  NumErrors += test<KName<Name, class Device>, false, T, BinaryOperation>(
       Q, Identity, Init, WGSize, NWItems, usm::alloc::device);
 }
 
 int main() {
   queue Q;
+  printDeviceInfo(Q);
+
   // fast atomics and fast reduce
-  testUSM<class AtomicReduce1, int, ONEAPI::plus<int>>(Q, 0, 99, 49, 5 * 49);
+  testUSM<class AtomicReduce1, int, std::plus<int>>(Q, 0, 99, 49, 5 * 49);
 
   // fast atomics
-  testUSM<class Atomic1, int, ONEAPI::bit_or<int>>(Q, 0, 0xff00ff00, 7, 7);
-  testUSM<class Atomic2, int, ONEAPI::bit_or<int>>(Q, 0, 0x7f007f00, 4, 32);
+  testUSM<class Atomic1, int, std::bit_or<>>(Q, 0, 0xff00ff00, 7, 7);
+  testUSM<class Atomic2, int, std::bit_or<>>(Q, 0, 0x7f007f00, 4, 32);
 
   // fast reduce
-  testUSM<class Reduce1, float, ONEAPI::minimum<float>>(
+  testUSM<class Reduce1, float, ext::oneapi::minimum<>>(
       Q, getMaximumFPValue<float>(), -100.0, 17, 17);
-  testUSM<class Reduce2, float, ONEAPI::maximum<float>>(
+  testUSM<class Reduce2, float, ext::oneapi::maximum<>>(
       Q, getMinimumFPValue<float>(), 100.0, 4, 32);
 
   // generic algorithm
-  testUSM<class Generic1, int, std::multiplies<int>>(Q, 1, 5, 7, 7);
+  testUSM<class Generic1, int, std::multiplies<>>(Q, 1, 5, 7, 7);
   testUSM<class Generic2, CustomVec<short>, CustomVecPlus<short>>(
       Q, CustomVec<short>(0), CustomVec<short>(77), 8, 8 * 3);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_usm_dw.cpp b/SYCL/Reduction/reduction_usm_dw.cpp
index e3c9b0d6a1..42f4f41a59 100644
--- a/SYCL/Reduction/reduction_usm_dw.cpp
+++ b/SYCL/Reduction/reduction_usm_dw.cpp
@@ -12,33 +12,30 @@
 // is created with property::reduction::initialize_to_identity.
 
 #include "reduction_utils.hpp"
-#include <CL/sycl.hpp>
-#include <cassert>
 
 using namespace cl::sycl;
 
-template <typename T1, typename T2> class KernelNameGroup;
+template <typename T1, typename T2> class KName;
+
+// Discard-write access to USM reductions is available only in SYCL2020.
+constexpr bool IsSYCL2020 = true;
 
 template <typename Name, typename T, class BinaryOperation>
-void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems,
-          usm::alloc AllocType) {
+int test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems,
+         usm::alloc AllocType) {
+  nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
+  printTestLabel<T, BinaryOperation>(IsSYCL2020, NDRange);
+
   auto Dev = Q.get_device();
-  if (AllocType == usm::alloc::shared &&
-      !Dev.get_info<info::device::usm_shared_allocations>())
-    return;
-  if (AllocType == usm::alloc::host &&
-      !Dev.get_info<info::device::usm_host_allocations>())
-    return;
-  if (AllocType == usm::alloc::device &&
-      !Dev.get_info<info::device::usm_device_allocations>())
-    return;
+  if (!Dev.has(getUSMAspect(AllocType)))
+    return 0;
 
   T *ReduVarPtr = (T *)malloc(sizeof(T), Dev, Q.get_context(), AllocType);
   if (ReduVarPtr == nullptr)
-    return;
+    return 0;
   if (AllocType == usm::alloc::device) {
     Q.submit([&](handler &CGH) {
-       CGH.single_task<KernelNameGroup<Name, class InitKernel>>(
+       CGH.single_task<KName<Name, class InitKernel>>(
            [=]() { *ReduVarPtr = Init; });
      }).wait();
   } else {
@@ -55,11 +52,9 @@ void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems,
   // Compute.
   Q.submit([&](handler &CGH) {
      auto In = InBuf.template get_access<access::mode::read>(CGH);
-     property_list PropList(property::reduction::initialize_to_identity{});
-     auto Redu = sycl::reduction(ReduVarPtr, Identity, BOp, PropList);
-     nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
-
-     CGH.parallel_for<KernelNameGroup<Name, class Test>>(
+     auto Redu = createReduction<IsSYCL2020, access::mode::discard_write>(
+         ReduVarPtr, Identity, BOp);
+     CGH.parallel_for<KName<Name, class Test>>(
          NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
            Sum.combine(In[NDIt.get_global_linear_id()]);
          });
@@ -71,54 +66,56 @@ void test(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems,
     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<KernelNameGroup<Name, class Check>>(
+       CGH.single_task<KName<Name, class Check>>(
            [=]() { OutAcc[0] = *ReduVarPtr; });
      }).wait();
     ComputedOut = (Buf.template get_access<access::mode::read>())[0];
   } else {
     ComputedOut = *ReduVarPtr;
   }
-  if (ComputedOut != CorrectOut) {
-    std::cerr << "NWItems = " << NWItems << ", WGSize = " << WGSize << "\n";
-    std::cerr << "Computed value: " << ComputedOut
-              << ", Expected value: " << CorrectOut
-              << ", AllocMode: " << static_cast<int>(AllocType) << "\n";
-    assert(0 && "Wrong value.");
-  }
+  std::string AllocStr =
+      "AllocMode=" + std::to_string(static_cast<int>(AllocType));
+  int Error = checkResults(Q, IsSYCL2020, BOp, NDRange, ComputedOut, CorrectOut,
+                           AllocStr);
 
   free(ReduVarPtr, Q.get_context());
+  return Error;
 }
 
+int NumErrors = 0;
+
 template <typename Name, typename T, class BinaryOperation>
 void testUSM(queue &Q, T Identity, T Init, size_t WGSize, size_t NWItems) {
-  test<KernelNameGroup<Name, class Shared2020>, T, BinaryOperation>(
+  NumErrors += test<KName<Name, class Shared2020>, T, BinaryOperation>(
       Q, Identity, Init, WGSize, NWItems, usm::alloc::shared);
-  test<KernelNameGroup<Name, class Host2020>, T, BinaryOperation>(
+  NumErrors += test<KName<Name, class Host2020>, T, BinaryOperation>(
       Q, Identity, Init, WGSize, NWItems, usm::alloc::host);
-  test<KernelNameGroup<Name, class Device2020>, T, BinaryOperation>(
+  NumErrors += test<KName<Name, class Device2020>, T, BinaryOperation>(
       Q, Identity, Init, WGSize, NWItems, usm::alloc::device);
 }
 
 int main() {
   queue Q;
+  printDeviceInfo(Q);
+
   // fast atomics and fast reduce
-  testUSM<class AtomicReduce1, int, ONEAPI::plus<int>>(Q, 0, 99, 49, 5 * 49);
+  testUSM<class AtomicReduce1, int, std::plus<int>>(Q, 0, 99, 49, 5 * 49);
 
   // fast atomics
-  testUSM<class Atomic1, int, ONEAPI::bit_or<int>>(Q, 0, 0xff00ff00, 7, 7);
-  testUSM<class Atomic2, int, ONEAPI::bit_or<int>>(Q, 0, 0x7f007f00, 4, 32);
+  testUSM<class Atomic1, int, std::bit_or<>>(Q, 0, 0xff00ff00, 7, 7);
+  testUSM<class Atomic2, int, std::bit_or<>>(Q, 0, 0x7f007f00, 4, 32);
 
   // fast reduce
-  testUSM<class Reduce1, float, ONEAPI::minimum<float>>(
+  testUSM<class Reduce1, float, ext::oneapi::minimum<>>(
       Q, getMaximumFPValue<float>(), -100.0, 17, 17);
-  testUSM<class Reduce2, float, ONEAPI::maximum<float>>(
+  testUSM<class Reduce2, float, ext::oneapi::maximum<>>(
       Q, getMinimumFPValue<float>(), 100.0, 4, 32);
 
   // generic algorithm
-  testUSM<class Generic1, int, std::multiplies<int>>(Q, 1, 5, 7, 7);
+  testUSM<class Generic1, int, std::multiplies<>>(Q, 1, 5, 7, 7);
   testUSM<class Generic2, CustomVec<short>, CustomVecPlus<short>>(
       Q, CustomVec<short>(0), CustomVec<short>(77), 8, 8 * 3);
 
-  std::cout << "Test passed\n";
-  return 0;
+  printFinalStatus(NumErrors);
+  return NumErrors;
 }
diff --git a/SYCL/Reduction/reduction_utils.hpp b/SYCL/Reduction/reduction_utils.hpp
index e51030f1ea..6303fe1d46 100644
--- a/SYCL/Reduction/reduction_utils.hpp
+++ b/SYCL/Reduction/reduction_utils.hpp
@@ -11,8 +11,9 @@ void initInputData(buffer<T, 1> &InBuf, T &ExpectedOut, T Identity,
   size_t N = Range.size();
   auto In = InBuf.template get_access<access::mode::write>();
   for (int I = 0; I < N; ++I) {
-    if (std::is_same<BinaryOperation, std::multiplies<T>>::value)
-      In[I] = 1 + (((I % 37) == 0) ? 1 : 0);
+    if (std::is_same_v<BinaryOperation, std::multiplies<T>> ||
+        std::is_same_v<BinaryOperation, std::multiplies<>>)
+      In[I] = 1.1 + (((I % 11) == 0) ? 1 : 0);
     else
       In[I] = ((I + 1) % 5) + 1.1;
     ExpectedOut = BOp(ExpectedOut, In[I]);
@@ -28,8 +29,9 @@ void initInputData(buffer<T, 2> &InBuf, T &ExpectedOut, T Identity,
   auto In = InBuf.template get_access<access::mode::write>();
   for (int J = 0; J < Range[0]; ++J) {
     for (int I = 0; I < Range[1]; ++I) {
-      if (std::is_same<BinaryOperation, std::multiplies<T>>::value)
-        In[J][I] = 1 + ((((I * 2 + J * 3) % 37) == 0) ? 1 : 0);
+      if (std::is_same_v<BinaryOperation, std::multiplies<T>> ||
+          std::is_same_v<BinaryOperation, std::multiplies<>>)
+        In[J][I] = 1.1 + ((((I + J * 3) % 11) == 0) ? 1 : 0);
       else
         In[J][I] = ((I + 1 + J) % 5) + 1.1;
       ExpectedOut = BOp(ExpectedOut, In[J][I]);
@@ -47,8 +49,9 @@ void initInputData(buffer<T, 3> &InBuf, T &ExpectedOut, T Identity,
   for (int K = 0; K < Range[0]; ++K) {
     for (int J = 0; J < Range[1]; ++J) {
       for (int I = 0; I < Range[2]; ++I) {
-        if (std::is_same<BinaryOperation, std::multiplies<T>>::value)
-          In[K][J][I] = 1 + ((((I * 2 + J * 3 + K) % 37) == 0) ? 1 : 0);
+        if (std::is_same_v<BinaryOperation, std::multiplies<T>> ||
+            std::is_same_v<BinaryOperation, std::multiplies<>>)
+          In[K][J][I] = 1.1 + ((((I + J * 3 + K) % 11) == 0) ? 1 : 0);
         else
           In[K][J][I] = ((I + 1 + J + K * 3) % 5) + 1.1;
         ExpectedOut = BOp(ExpectedOut, In[K][J][I]);
@@ -75,6 +78,23 @@ bool operator==(const CustomVec<T> &A, const CustomVec<T> &B) {
   return A.X == B.X && A.Y == B.Y;
 }
 template <typename T>
+bool operator<(const CustomVec<T> &A, const CustomVec<T> &B) {
+  return A.X < B.X && A.Y < B.Y;
+}
+template <typename T>
+CustomVec<T> operator/(const CustomVec<T> &A, const CustomVec<T> &B) {
+  return {A.X / B.X && A.Y / B.Y};
+}
+template <typename T>
+CustomVec<T> operator-(const CustomVec<T> &A, const CustomVec<T> &B) {
+  return {A.X - B.X && A.Y - B.Y};
+}
+namespace std {
+template <typename T> CustomVec<T> abs(const CustomVec<T> &A) {
+  return {std::abs(A.X), std::abs(A.Y)};
+}
+} // namespace std
+template <typename T>
 std::ostream &operator<<(std::ostream &OS, const CustomVec<T> &V) {
   return OS << "(" << V.X << ", " << V.Y << ")";
 }
@@ -105,14 +125,142 @@ template <access::mode Mode> property_list getPropertyList() {
 }
 
 void printDeviceInfo(queue &Q, bool ToCERR = false) {
+  static int IsErrDeviceInfoPrinted = 0;
+  if (IsErrDeviceInfoPrinted >= 2)
+    return;
+  IsErrDeviceInfoPrinted++;
+
   device D = Q.get_device();
   auto Name = D.get_info<sycl::info::device::name>();
   size_t MaxWGSize = D.get_info<info::device::max_work_group_size>();
   size_t LocalMemSize = D.get_info<info::device::local_mem_size>();
-  if (ToCERR)
-    std::cout << "Device: " << Name << ", MaxWGSize: " << MaxWGSize
-              << ", LocalMemSize: " << LocalMemSize << std::endl;
+  std::ostream &OS = ToCERR ? std::cerr : std::cout;
+  OS << "Device: " << Name << ", MaxWGSize: " << MaxWGSize
+     << ", LocalMemSize: " << LocalMemSize
+     << ", Driver: " << D.get_info<info::device::driver_version>() << std::endl;
+}
+
+template <int Dims>
+std::ostream &operator<<(std::ostream &OS, const range<Dims> &Range) {
+  OS << "{" << Range[0];
+  if constexpr (Dims > 1)
+    OS << ", " << Range[1];
+  if constexpr (Dims > 2)
+    OS << ", " << Range[2];
+  OS << "}";
+  return OS;
+}
+
+template <int Dims>
+std::ostream &operator<<(std::ostream &OS, const nd_range<Dims> &Range) {
+  OS << "{" << Range.get_global_range() << ", " << Range.get_local_range()
+     << "}";
+  return OS;
+}
+
+template <typename T, typename BinaryOperation, typename RangeT>
+void printTestLabel(bool IsSYCL2020, const RangeT &Range, bool ToCERR = false) {
+  std::ostream &OS = ToCERR ? std::cerr : std::cout;
+  std::string Mode = IsSYCL2020 ? "SYCL2020" : "ONEAPI  ";
+  OS << (ToCERR ? "Error" : "Start") << ": Mode=" << Mode
+     << ", T=" << typeid(T).name() << ", BOp=" << typeid(BinaryOperation).name()
+     << ", Range=" << Range;
+}
+
+template <typename BOp, typename T> constexpr bool isPreciseResultFP() {
+  return (std::is_floating_point_v<T> || std::is_same_v<T, half>)&&(
+      std::is_same_v<ext::oneapi::minimum<>, BOp> ||
+      std::is_same_v<ext::oneapi::minimum<T>, BOp> ||
+      std::is_same_v<ext::oneapi::maximum<>, BOp> ||
+      std::is_same_v<ext::oneapi::maximum<T>, BOp>);
+}
+
+template <typename BinaryOperation, typename T, typename RangeT>
+int checkResults(queue &Q, bool IsSYCL2020, BinaryOperation,
+                 const RangeT &Range, const T &ComputedRes, const T &CorrectRes,
+                 std::string AddInfo = "") {
+  std::string ErrorStr;
+  bool Passed;
+
+  if constexpr (std::is_floating_point_v<T> || std::is_same_v<T, half>) {
+    // It is a pretty simple and naive FP diff check here, which though
+    // should work reasonably well for most of cases handled in reduction
+    // tests.
+    T MaxDiff = std::numeric_limits<T>::epsilon() * std::fabs(CorrectRes);
+    if constexpr (std::is_same_v<RangeT, range<1>> ||
+                  std::is_same_v<RangeT, range<2>> ||
+                  std::is_same_v<RangeT, range<3>>)
+      MaxDiff *= Range.size();
+    else
+      MaxDiff *= Range.get_global_range().size();
+
+    if (isPreciseResultFP<BinaryOperation, T>())
+      MaxDiff = 0;
+    T Diff = std::abs(CorrectRes - ComputedRes);
+    ErrorStr = ", Diff=" + std::to_string(Diff) +
+               ", MaxDiff=" + std::to_string(MaxDiff);
+    Passed = Diff <= MaxDiff;
+  } else {
+    Passed = ComputedRes == CorrectRes;
+  }
+
+  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;
+  }
+  return Passed ? 0 : 1;
+}
+
+void printFinalStatus(int NumErrors) {
+  if (NumErrors == 0)
+    std::cout << "Test passed" << std::endl;
   else
-    std::cerr << "Device: " << Name << ", MaxWGSize: " << MaxWGSize
-              << ", LocalMemSize: " << LocalMemSize << std::endl;
+    std::cerr << NumErrors << " test-cases failed" << std::endl;
+}
+
+template <bool IsSYCL2020, access::mode AccMode = access::mode::read_write,
+          typename T, typename BinaryOperation>
+auto createReduction(T *USMPtr, T Identity, BinaryOperation BOp) {
+  if constexpr (IsSYCL2020)
+    return sycl::reduction(USMPtr, Identity, BOp, getPropertyList<AccMode>());
+  else
+    return ext::oneapi::reduction(USMPtr, Identity, BOp);
+}
+
+template <bool IsSYCL2020, access::mode AccMode = access::mode::read_write,
+          typename T, typename BinaryOperation>
+auto createReduction(T *USMPtr, BinaryOperation BOp) {
+  if constexpr (IsSYCL2020)
+    return sycl::reduction(USMPtr, BOp, getPropertyList<AccMode>());
+  else
+    return ext::oneapi::reduction(USMPtr, BOp);
+}
+
+template <bool IsSYCL2020, access::mode AccMode, int AccDim = 1, typename T,
+          typename BinaryOperation, typename BufferT>
+auto createReduction(BufferT ReduBuf, handler &CGH, T Identity,
+                     BinaryOperation BOp) {
+  if constexpr (IsSYCL2020) {
+    property_list PropList = getPropertyList<AccMode>();
+    return sycl::reduction(ReduBuf, CGH, Identity, BOp, PropList);
+  } else {
+    accessor<T, AccDim, AccMode, access::target::global_buffer> Out(ReduBuf,
+                                                                    CGH);
+    return ext::oneapi::reduction(Out, Identity, BOp);
+  }
+}
+
+aspect getUSMAspect(usm::alloc Alloc) {
+  if (Alloc == sycl::usm::alloc::host)
+    return aspect::usm_host_allocations;
+
+  if (Alloc == sycl::usm::alloc::device)
+    return aspect::usm_device_allocations;
+
+  assert(Alloc == usm::alloc::shared && "Unknown USM allocation type");
+  return aspect::usm_shared_allocations;
 }