From 82dbc8ac1ba1204702de56e7b98f4c1092aee182 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Tue, 8 Nov 2022 12:39:25 -0800
Subject: [PATCH 1/3] [NFCI][SYCL] Refactor reductions implementations

The purpose of this change is

1) Make it clear that sycl::range version always delegates to some
   sycl::nd_range implementation.

2) Flatten all existing implementations. That required splitting
   reduction_parallel_for_basic_impl depending on
   Reduction::has_fast_reduce. However, I also inlined its helper
   routines (that now had unambigious branches for "if
   constexpr (Reduction::*)") so the ratio between duplicate/unique code
   isn't bad at all.

3) Make a unique dispatching entry and have all the implementations
   provide the same interface. I plan to use it in unit-tests to bypass
   the dispatch and test all the implementation directly (when
   applicable based on nd_range/BinOp/HW/etc.).

I'd say 90% of the change is straightforward code movement with the
exceptions of

  - Inlining described above
  - Simplifying __sycl_reduction_kernel helper enabled by this change
  - Factoring out forward declarations in handler.hpp into a separate
    reduction_fwd.hpp
  - Rewriting dispatcher routine - caused by changes in the interfaces,
    not logic updates.
---
 sycl/include/sycl/handler.hpp                 |   35 +-
 sycl/include/sycl/reduction.hpp               | 1516 ++++++++---------
 sycl/include/sycl/reduction_fwd.hpp           |   67 +
 .../properties_kernel_device_has.cpp          |   24 +-
 .../properties_kernel_sub_group_size.cpp      |   72 +-
 .../properties_kernel_work_group_size.cpp     |   72 +-
 ...properties_kernel_work_group_size_hint.cpp |   72 +-
 7 files changed, 942 insertions(+), 916 deletions(-)
 create mode 100644 sycl/include/sycl/reduction_fwd.hpp

diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 971c91de9c5a8..c03db5fd23e9a 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -30,6 +30,7 @@
 #include <sycl/nd_item.hpp>
 #include <sycl/nd_range.hpp>
 #include <sycl/property_list.hpp>
+#include <sycl/reduction_fwd.hpp>
 #include <sycl/sampler.hpp>
 #include <sycl/stl.hpp>
 
@@ -268,43 +269,9 @@ class RoundedRangeKernelWithKH {
   KernelType KernelFunc;
 };
 
-template <typename T, class BinaryOperation, int Dims, size_t Extent,
-          typename RedOutVar>
-class reduction_impl_algo;
-
 using sycl::detail::enable_if_t;
 using sycl::detail::queue_impl;
 
-// Reductions implementation need access to private members of handler. Those
-// are limited to those below.
-namespace reduction {
-inline void finalizeHandler(handler &CGH);
-template <class FunctorTy> void withAuxHandler(handler &CGH, FunctorTy Func);
-} // namespace reduction
-
-template <typename KernelName, int Dims, typename PropertiesT,
-          typename KernelType, typename Reduction>
-void reduction_parallel_for(handler &CGH,
-                            std::shared_ptr<detail::queue_impl> Queue,
-                            range<Dims> Range, PropertiesT Properties,
-                            Reduction Redu, KernelType KernelFunc);
-
-template <typename KernelName, int Dims, typename PropertiesT,
-          typename KernelType, typename Reduction>
-void reduction_parallel_for(handler &CGH,
-                            std::shared_ptr<detail::queue_impl> Queue,
-                            nd_range<Dims> Range, PropertiesT Properties,
-                            Reduction Redu, KernelType KernelFunc);
-
-template <typename KernelName, int Dims, typename PropertiesT,
-          typename... RestT>
-void reduction_parallel_for(handler &CGH,
-                            std::shared_ptr<detail::queue_impl> Queue,
-                            nd_range<Dims> Range, PropertiesT Properties,
-                            RestT... Rest);
-
-template <typename T> struct IsReduction;
-template <typename FirstT, typename... RestT> struct AreAllButLastReductions;
 } // namespace detail
 
 /// Command group handler class.
diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index 6eeb263aa0ca1..14bd060d6e54b 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -18,6 +18,7 @@
 #include <sycl/kernel.hpp>
 #include <sycl/known_identity.hpp>
 #include <sycl/properties/reduction_properties.hpp>
+#include <sycl/reduction_fwd.hpp>
 #include <sycl/usm.hpp>
 
 #include <tuple>
@@ -317,7 +318,7 @@ template <class Reducer> class combiner {
         ReduVarPtr, [](auto Ref, auto Val) { return Ref.fetch_max(Val); });
   }
 };
-}
+} // namespace detail
 
 /// Specialization of the generic class 'reducer'. It is used for reductions
 /// of those types and operations for which the identity value is not known.
@@ -571,8 +572,7 @@ class reduction_impl_algo : public reduction_impl_common<T, BinaryOperation> {
     }
   }
 
-  template <class _T = T>
-  auto &getTempBuffer(size_t Size, handler &CGH) {
+  template <class _T = T> auto &getTempBuffer(size_t Size, handler &CGH) {
     auto Buffer = std::make_shared<buffer<_T, 1>>(range<1>(Size));
     CGH.addReduction(Buffer);
     return *Buffer;
@@ -830,711 +830,785 @@ reduSaveFinalResultToUserMem(handler &CGH, Reduction &Redu) {
   });
 }
 
+namespace reduction {
+template <typename KernelName, strategy S, class... Ts> struct MainKrn;
+template <typename KernelName, strategy S, class... Ts> struct AuxKrn;
+} // namespace reduction
+
 /// A helper to pass undefined (sycl::detail::auto_name) names unmodified. We
 /// must do that to avoid name collisions.
-template <template <typename...> class Namer, class KernelName, class... Ts>
+template <template <typename, reduction::strategy, typename...> class MainOrAux,
+          class KernelName, reduction::strategy Strategy, class... Ts>
 using __sycl_reduction_kernel =
     std::conditional_t<std::is_same<KernelName, auto_name>::value, auto_name,
-                       Namer<KernelName, Ts...>>;
+                       MainOrAux<KernelName, Strategy, Ts...>>;
 
-namespace reduction {
-namespace main_krn {
-template <class KernelName> struct RangeFastAtomics;
-} // namespace main_krn
-} // namespace reduction
-template <typename KernelName, typename KernelType, typename PropertiesT,
-          class Reduction>
-void reduCGFuncForRangeFastAtomics(handler &CGH, KernelType KernelFunc,
-                                   const nd_range<1> &NDRange,
-                                   PropertiesT Properties, Reduction &Redu) {
-  size_t NElements = Reduction::num_elements;
-  auto Out = Redu.getReadWriteAccessorToInitializedMem(CGH);
-  local_accessor<typename Reduction::result_type, 1> GroupSum{NElements, CGH};
-  using Name = __sycl_reduction_kernel<reduction::main_krn::RangeFastAtomics,
-                                       KernelName>;
-  CGH.parallel_for<Name>(NDRange, Properties, [=](nd_item<1> NDId) {
-    // Call user's functions. Reducer.MValue gets initialized there.
-    typename Reduction::reducer_type Reducer;
-    KernelFunc(NDId, Reducer);
-
-    // Work-group cooperates to initialize multiple reduction variables
-    auto LID = NDId.get_local_id(0);
-    for (size_t E = LID; E < NElements; E += NDId.get_local_range(0)) {
-      GroupSum[E] = Reducer.getIdentity();
-    }
-    workGroupBarrier();
+// Implementations.
 
-    // Each work-item has its own reducer to combine
-    Reducer.template atomic_combine<access::address_space::local_space>(
-        &GroupSum[0]);
+template <reduction::strategy> struct NDRangeReduction;
 
-    // Single work-item performs finalization for entire work-group
-    // TODO: Opportunity to parallelize across elements
-    workGroupBarrier();
-    if (LID == 0) {
-      for (size_t E = 0; E < NElements; ++E) {
-        Reducer.getElement(E) = GroupSum[E];
+template <>
+struct NDRangeReduction<reduction::strategy::local_atomic_and_atomic_cross_wg> {
+  template <typename KernelName, int Dims, typename PropertiesT,
+            typename KernelType, typename Reduction>
+  static void run(handler &CGH, std::shared_ptr<detail::queue_impl> &Queue,
+                  nd_range<Dims> NDRange, PropertiesT &Properties,
+                  Reduction &Redu, KernelType &KernelFunc) {
+    size_t NElements = Reduction::num_elements;
+    auto Out = Redu.getReadWriteAccessorToInitializedMem(CGH);
+    local_accessor<typename Reduction::result_type, 1> GroupSum{NElements, CGH};
+
+    using Name = __sycl_reduction_kernel<
+        reduction::MainKrn, KernelName,
+        reduction::strategy::local_atomic_and_atomic_cross_wg>;
+
+    CGH.parallel_for<Name>(NDRange, Properties, [=](nd_item<1> NDId) {
+      // Call user's functions. Reducer.MValue gets initialized there.
+      typename Reduction::reducer_type Reducer;
+      KernelFunc(NDId, Reducer);
+
+      // Work-group cooperates to initialize multiple reduction variables
+      auto LID = NDId.get_local_id(0);
+      for (size_t E = LID; E < NElements; E += NDId.get_local_range(0)) {
+        GroupSum[E] = Reducer.getIdentity();
       }
-      Reducer.template atomic_combine(&Out[0]);
-    }
-  });
+      workGroupBarrier();
 
-  if (Reduction::is_usm || Redu.initializeToIdentity())
-    reduction::withAuxHandler(CGH, [&](handler &CopyHandler) {
-      reduSaveFinalResultToUserMem<KernelName>(CopyHandler, Redu);
+      // Each work-item has its own reducer to combine
+      Reducer.template atomic_combine<access::address_space::local_space>(
+          &GroupSum[0]);
+
+      // Single work-item performs finalization for entire work-group
+      // TODO: Opportunity to parallelize across elements
+      workGroupBarrier();
+      if (LID == 0) {
+        for (size_t E = 0; E < NElements; ++E) {
+          Reducer.getElement(E) = GroupSum[E];
+        }
+        Reducer.template atomic_combine(&Out[0]);
+      }
     });
-}
 
-namespace reduction {
-namespace main_krn {
-template <class KernelName, class NWorkGroupsFinished> struct RangeFastReduce;
-} // namespace main_krn
-} // namespace reduction
-template <typename KernelName, typename KernelType, typename PropertiesT,
-          class Reduction>
-void reduCGFuncForRangeFastReduce(handler &CGH, KernelType KernelFunc,
-                                  const nd_range<1> &NDRange,
-                                  PropertiesT Properties, Reduction &Redu) {
-  size_t NElements = Reduction::num_elements;
-  size_t WGSize = NDRange.get_local_range().size();
-  size_t NWorkGroups = NDRange.get_group_range().size();
+    if (Reduction::is_usm || Redu.initializeToIdentity())
+      reduction::withAuxHandler(CGH, [&](handler &CopyHandler) {
+        reduSaveFinalResultToUserMem<KernelName>(CopyHandler, Redu);
+      });
+  }
+};
+
+template <>
+struct NDRangeReduction<
+    reduction::strategy::group_reduce_and_last_wg_detection> {
+  template <typename KernelName, int Dims, typename PropertiesT,
+            typename KernelType, typename Reduction>
+  static void run(handler &CGH, std::shared_ptr<detail::queue_impl> &Queue,
+                  nd_range<Dims> NDRange, PropertiesT &Properties,
+                  Reduction &Redu, KernelType &KernelFunc) {
+    size_t NElements = Reduction::num_elements;
+    size_t WGSize = NDRange.get_local_range().size();
+    size_t NWorkGroups = NDRange.get_group_range().size();
+
+    auto &Out = Redu.getUserRedVar();
+    if constexpr (!Reduction::is_usm)
+      associateWithHandler(CGH, &Out, access::target::device);
+
+    auto &PartialSumsBuf = Redu.getTempBuffer(NWorkGroups * NElements, CGH);
+    accessor PartialSums(PartialSumsBuf, CGH, sycl::read_write, sycl::no_init);
+
+    bool IsUpdateOfUserVar = !Reduction::is_usm && !Redu.initializeToIdentity();
+    auto Rest = [&](auto NWorkGroupsFinished) {
+      local_accessor<int, 1> DoReducePartialSumsInLastWG{1, CGH};
+
+      using Name = __sycl_reduction_kernel<
+          reduction::MainKrn, KernelName,
+          reduction::strategy::group_reduce_and_last_wg_detection,
+          decltype(NWorkGroupsFinished)>;
+
+      CGH.parallel_for<Name>(NDRange, Properties, [=](nd_item<1> NDId) {
+        // Call user's functions. Reducer.MValue gets initialized there.
+        typename Reduction::reducer_type Reducer;
+        KernelFunc(NDId, Reducer);
+
+        typename Reduction::binary_operation BOp;
+        auto Group = NDId.get_group();
+
+        // If there are multiple values, reduce each separately
+        // reduce_over_group is only defined for each T, not for span<T, ...>
+        size_t LID = NDId.get_local_id(0);
+        for (int E = 0; E < NElements; ++E) {
+          auto &RedElem = Reducer.getElement(E);
+          RedElem = reduce_over_group(Group, RedElem, BOp);
+          if (LID == 0) {
+            if (NWorkGroups == 1) {
+              // Can avoid using partial sum and write the final result
+              // immediately.
+              if (IsUpdateOfUserVar)
+                RedElem = BOp(RedElem, Out[E]);
+              Out[E] = RedElem;
+            } else {
+              PartialSums[NDId.get_group_linear_id() * NElements + E] =
+                  Reducer.getElement(E);
+            }
+          }
+        }
 
-  auto &Out = Redu.getUserRedVar();
-  if constexpr (!Reduction::is_usm)
-    associateWithHandler(CGH, &Out, access::target::device);
+        if (NWorkGroups == 1)
+          // We're done.
+          return;
 
-  auto &PartialSumsBuf = Redu.getTempBuffer(NWorkGroups * NElements, CGH);
-  accessor PartialSums(PartialSumsBuf, CGH, sycl::read_write, sycl::no_init);
+        // Signal this work-group has finished after all values are reduced
+        if (LID == 0) {
+          auto NFinished =
+              sycl::atomic_ref<int, memory_order::relaxed, memory_scope::device,
+                               access::address_space::global_space>(
+                  NWorkGroupsFinished[0]);
+          DoReducePartialSumsInLastWG[0] = ++NFinished == NWorkGroups;
+        }
 
-  bool IsUpdateOfUserVar = !Reduction::is_usm && !Redu.initializeToIdentity();
-  auto Rest = [&](auto NWorkGroupsFinished) {
+        workGroupBarrier();
+        if (DoReducePartialSumsInLastWG[0]) {
+          // Reduce each result separately
+          // TODO: Opportunity to parallelize across elements.
+          for (int E = 0; E < NElements; ++E) {
+            auto LocalSum = Reducer.getIdentity();
+            for (size_t I = LID; I < NWorkGroups; I += WGSize)
+              LocalSum = BOp(LocalSum, PartialSums[I * NElements + E]);
+            auto Result = reduce_over_group(Group, LocalSum, BOp);
+
+            if (LID == 0) {
+              if (IsUpdateOfUserVar)
+                Result = BOp(Result, Out[E]);
+              Out[E] = Result;
+            }
+          }
+        }
+      });
+    };
+
+    auto device = getDeviceFromHandler(CGH);
+    // Integrated/discrete GPUs have different faster path. For discrete GPUs
+    // fast path requires USM device allocations though, so check for that as
+    // well.
+    if (device.get_info<info::device::host_unified_memory>() ||
+        !device.has(aspect::usm_device_allocations))
+      Rest(Redu.getReadWriteAccessorToInitializedGroupsCounter(CGH));
+    else
+      Rest(Redu.getGroupsCounterAccDiscrete(CGH));
+  }
+};
+
+template <> struct NDRangeReduction<reduction::strategy::range_basic> {
+  template <typename KernelName, int Dims, typename PropertiesT,
+            typename KernelType, typename Reduction>
+  static void run(handler &CGH, std::shared_ptr<detail::queue_impl> &Queue,
+                  nd_range<Dims> NDRange, PropertiesT &Properties,
+                  Reduction &Redu, KernelType &KernelFunc) {
+    size_t NElements = Reduction::num_elements;
+    size_t WGSize = NDRange.get_local_range().size();
+    size_t NWorkGroups = NDRange.get_group_range().size();
+
+    bool IsUpdateOfUserVar = !Reduction::is_usm && !Redu.initializeToIdentity();
+    auto PartialSums =
+        Redu.getWriteAccForPartialReds(NWorkGroups * NElements, CGH);
+    auto Out = (NWorkGroups == 1)
+                   ? PartialSums
+                   : Redu.getWriteAccForPartialReds(NElements, CGH);
+    local_accessor<typename Reduction::result_type, 1> LocalReds{WGSize + 1,
+                                                                 CGH};
+    auto NWorkGroupsFinished =
+        Redu.getReadWriteAccessorToInitializedGroupsCounter(CGH);
     local_accessor<int, 1> DoReducePartialSumsInLastWG{1, CGH};
 
-    using Name =
-        __sycl_reduction_kernel<reduction::main_krn::RangeFastReduce,
-                                KernelName, decltype(NWorkGroupsFinished)>;
+    auto Identity = Redu.getIdentity();
+    auto BOp = Redu.getBinaryOperation();
+
+    using Name = __sycl_reduction_kernel<reduction::MainKrn, KernelName,
+                                         reduction::strategy::range_basic>;
+
     CGH.parallel_for<Name>(NDRange, Properties, [=](nd_item<1> NDId) {
       // Call user's functions. Reducer.MValue gets initialized there.
-      typename Reduction::reducer_type Reducer;
+      typename Reduction::reducer_type Reducer(Identity, BOp);
       KernelFunc(NDId, Reducer);
 
-      typename Reduction::binary_operation BOp;
-      auto Group = NDId.get_group();
-
       // If there are multiple values, reduce each separately
-      // reduce_over_group is only defined for each T, not for span<T, ...>
-      size_t LID = NDId.get_local_id(0);
+      // This prevents local memory from scaling with elements
+      size_t LID = NDId.get_local_linear_id();
       for (int E = 0; E < NElements; ++E) {
-        auto &RedElem = Reducer.getElement(E);
-        RedElem = reduce_over_group(Group, RedElem, BOp);
+
+        // Copy the element to local memory to prepare it for tree-reduction.
+        LocalReds[LID] = Reducer.getElement(E);
+        if (LID == 0)
+          LocalReds[WGSize] = Identity;
+        workGroupBarrier();
+
+        // Tree-reduction: reduce the local array LocalReds[:] to LocalReds[0].
+        // LocalReds[WGSize] accumulates last/odd elements when the step
+        // of tree-reduction loop is not even.
+        size_t PrevStep = WGSize;
+        for (size_t CurStep = PrevStep >> 1; CurStep > 0; CurStep >>= 1) {
+          if (LID < CurStep)
+            LocalReds[LID] = BOp(LocalReds[LID], LocalReds[LID + CurStep]);
+          else if (LID == CurStep && (PrevStep & 0x1))
+            LocalReds[WGSize] = BOp(LocalReds[WGSize], LocalReds[PrevStep - 1]);
+          workGroupBarrier();
+          PrevStep = CurStep;
+        }
+
         if (LID == 0) {
-          if (NWorkGroups == 1) {
-            // Can avoid using partial sum and write the final result
-            // immediately.
-            if (IsUpdateOfUserVar)
-              RedElem = BOp(RedElem, Out[E]);
-            Out[E] = RedElem;
-          } else {
-            PartialSums[NDId.get_group_linear_id() * NElements + E] =
-                Reducer.getElement(E);
-          }
+          auto V = BOp(LocalReds[0], LocalReds[WGSize]);
+          if (NWorkGroups == 1 && IsUpdateOfUserVar)
+            V = BOp(V, Out[E]);
+          // if NWorkGroups == 1, then PartialsSum and Out point to same memory.
+          PartialSums[NDId.get_group_linear_id() * NElements + E] = V;
         }
       }
 
-      if (NWorkGroups == 1)
-        // We're done.
-        return;
-
       // Signal this work-group has finished after all values are reduced
       if (LID == 0) {
         auto NFinished =
             sycl::atomic_ref<int, memory_order::relaxed, memory_scope::device,
                              access::address_space::global_space>(
                 NWorkGroupsFinished[0]);
-        DoReducePartialSumsInLastWG[0] = ++NFinished == NWorkGroups;
+        DoReducePartialSumsInLastWG[0] =
+            ++NFinished == NWorkGroups && NWorkGroups > 1;
       }
 
       workGroupBarrier();
       if (DoReducePartialSumsInLastWG[0]) {
         // Reduce each result separately
-        // TODO: Opportunity to parallelize across elements.
+        // TODO: Opportunity to parallelize across elements
         for (int E = 0; E < NElements; ++E) {
-          auto LocalSum = Reducer.getIdentity();
+          auto LocalSum = Identity;
           for (size_t I = LID; I < NWorkGroups; I += WGSize)
             LocalSum = BOp(LocalSum, PartialSums[I * NElements + E]);
-          auto Result = reduce_over_group(Group, LocalSum, BOp);
 
+          LocalReds[LID] = LocalSum;
+          if (LID == 0)
+            LocalReds[WGSize] = Identity;
+          workGroupBarrier();
+
+          size_t PrevStep = WGSize;
+          for (size_t CurStep = PrevStep >> 1; CurStep > 0; CurStep >>= 1) {
+            if (LID < CurStep)
+              LocalReds[LID] = BOp(LocalReds[LID], LocalReds[LID + CurStep]);
+            else if (LID == CurStep && (PrevStep & 0x1))
+              LocalReds[WGSize] =
+                  BOp(LocalReds[WGSize], LocalReds[PrevStep - 1]);
+            workGroupBarrier();
+            PrevStep = CurStep;
+          }
           if (LID == 0) {
+            auto V = BOp(LocalReds[0], LocalReds[WGSize]);
             if (IsUpdateOfUserVar)
-              Result = BOp(Result, Out[E]);
-            Out[E] = Result;
+              V = BOp(V, Out[E]);
+            Out[E] = V;
           }
         }
       }
     });
-  };
 
-  auto device = getDeviceFromHandler(CGH);
-  // Integrated/discrete GPUs have different faster path. For discrete GPUs fast
-  // path requires USM device allocations though, so check for that as well.
-  if (device.get_info<info::device::host_unified_memory>() ||
-      !device.has(aspect::usm_device_allocations))
-    Rest(Redu.getReadWriteAccessorToInitializedGroupsCounter(CGH));
-  else
-    Rest(Redu.getGroupsCounterAccDiscrete(CGH));
-}
+    if (Reduction::is_usm)
+      reduction::withAuxHandler(CGH, [&](handler &CopyHandler) {
+        reduSaveFinalResultToUserMem<KernelName>(CopyHandler, Redu);
+      });
+  }
+};
 
-namespace reduction {
-namespace main_krn {
-template <class KernelName> struct RangeBasic;
-} // namespace main_krn
-} // namespace reduction
-template <typename KernelName, typename KernelType, typename PropertiesT,
-          class Reduction>
-void reduCGFuncForRangeBasic(handler &CGH, KernelType KernelFunc,
-                             const nd_range<1> &NDRange, PropertiesT Properties,
-                             Reduction &Redu) {
-  size_t NElements = Reduction::num_elements;
-  size_t WGSize = NDRange.get_local_range().size();
-  size_t NWorkGroups = NDRange.get_group_range().size();
-
-  bool IsUpdateOfUserVar = !Reduction::is_usm && !Redu.initializeToIdentity();
-  auto PartialSums =
-      Redu.getWriteAccForPartialReds(NWorkGroups * NElements, CGH);
-  auto Out = (NWorkGroups == 1)
-                 ? PartialSums
-                 : Redu.getWriteAccForPartialReds(NElements, CGH);
-  local_accessor<typename Reduction::result_type, 1> LocalReds{WGSize + 1, CGH};
-  auto NWorkGroupsFinished =
-      Redu.getReadWriteAccessorToInitializedGroupsCounter(CGH);
-  local_accessor<int, 1> DoReducePartialSumsInLastWG{1, CGH};
-
-  auto Identity = Redu.getIdentity();
-  auto BOp = Redu.getBinaryOperation();
-  using Name =
-      __sycl_reduction_kernel<reduction::main_krn::RangeBasic, KernelName>;
-  CGH.parallel_for<Name>(NDRange, Properties, [=](nd_item<1> NDId) {
-    // Call user's functions. Reducer.MValue gets initialized there.
-    typename Reduction::reducer_type Reducer(Identity, BOp);
-    KernelFunc(NDId, Reducer);
-
-    // If there are multiple values, reduce each separately
-    // This prevents local memory from scaling with elements
-    size_t LID = NDId.get_local_linear_id();
-    for (int E = 0; E < NElements; ++E) {
-
-      // Copy the element to local memory to prepare it for tree-reduction.
-      LocalReds[LID] = Reducer.getElement(E);
-      if (LID == 0)
-        LocalReds[WGSize] = Identity;
-      workGroupBarrier();
+template <>
+struct NDRangeReduction<reduction::strategy::group_reduce_and_atomic_cross_wg> {
+  template <typename KernelName, int Dims, typename PropertiesT,
+            typename KernelType, typename Reduction>
+  static void run(handler &CGH, std::shared_ptr<detail::queue_impl> &Queue,
+                  nd_range<Dims> NDRange, PropertiesT &Properties,
+                  Reduction &Redu, KernelType &KernelFunc) {
+    auto Out = Redu.getReadWriteAccessorToInitializedMem(CGH);
+    size_t NElements = Reduction::num_elements;
+
+    using Name = __sycl_reduction_kernel<
+        reduction::MainKrn, KernelName,
+        reduction::strategy::group_reduce_and_atomic_cross_wg>;
+
+    CGH.parallel_for<Name>(NDRange, Properties, [=](nd_item<Dims> NDIt) {
+      // Call user's function. Reducer.MValue gets initialized there.
+      typename Reduction::reducer_type Reducer;
+      KernelFunc(NDIt, Reducer);
 
-      // Tree-reduction: reduce the local array LocalReds[:] to LocalReds[0].
-      // LocalReds[WGSize] accumulates last/odd elements when the step
-      // of tree-reduction loop is not even.
-      size_t PrevStep = WGSize;
-      for (size_t CurStep = PrevStep >> 1; CurStep > 0; CurStep >>= 1) {
-        if (LID < CurStep)
-          LocalReds[LID] = BOp(LocalReds[LID], LocalReds[LID + CurStep]);
-        else if (LID == CurStep && (PrevStep & 0x1))
-          LocalReds[WGSize] = BOp(LocalReds[WGSize], LocalReds[PrevStep - 1]);
-        workGroupBarrier();
-        PrevStep = CurStep;
+      typename Reduction::binary_operation BOp;
+      for (int E = 0; E < NElements; ++E) {
+        Reducer.getElement(E) =
+            reduce_over_group(NDIt.get_group(), Reducer.getElement(E), BOp);
       }
+      if (NDIt.get_local_linear_id() == 0)
+        Reducer.atomic_combine(&Out[0]);
+    });
 
-      if (LID == 0) {
-        auto V = BOp(LocalReds[0], LocalReds[WGSize]);
-        if (NWorkGroups == 1 && IsUpdateOfUserVar)
-          V = BOp(V, Out[E]);
-        // if NWorkGroups == 1, then PartialsSum and Out point to same memory.
-        PartialSums[NDId.get_group_linear_id() * NElements + E] = V;
-      }
+    if (Reduction::is_usm || Redu.initializeToIdentity()) {
+      reduction::withAuxHandler(CGH, [&](handler &CopyHandler) {
+        reduSaveFinalResultToUserMem<KernelName>(CopyHandler, Redu);
+      });
     }
+  }
+};
 
-    // Signal this work-group has finished after all values are reduced
-    if (LID == 0) {
-      auto NFinished =
-          sycl::atomic_ref<int, memory_order::relaxed, memory_scope::device,
-                           access::address_space::global_space>(
-              NWorkGroupsFinished[0]);
-      DoReducePartialSumsInLastWG[0] =
-          ++NFinished == NWorkGroups && NWorkGroups > 1;
-    }
+template <>
+struct NDRangeReduction<
+    reduction::strategy::local_mem_tree_and_atomic_cross_wg> {
+  template <typename KernelName, int Dims, typename PropertiesT,
+            typename KernelType, typename Reduction>
+  static void run(handler &CGH, std::shared_ptr<detail::queue_impl> &Queue,
+                  nd_range<Dims> NDRange, PropertiesT &Properties,
+                  Reduction &Redu, KernelType &KernelFunc) {
+    auto Out = Redu.getReadWriteAccessorToInitializedMem(CGH);
+    size_t NElements = Reduction::num_elements;
+    size_t WGSize = NDRange.get_local_range().size();
+    bool IsPow2WG = (WGSize & (WGSize - 1)) == 0;
+
+    // Use local memory to reduce elements in work-groups into zero-th element.
+    // If WGSize is not power of two, then WGSize+1 elements are allocated.
+    // The additional last element is used to catch reduce elements that could
+    // otherwise be lost in the tree-reduction algorithm used in the kernel.
+    size_t NLocalElements = WGSize + (IsPow2WG ? 0 : 1);
+    local_accessor<typename Reduction::result_type, 1> LocalReds{NLocalElements,
+                                                                 CGH};
+
+    using Name = __sycl_reduction_kernel<
+        reduction::MainKrn, KernelName,
+        reduction::strategy::local_mem_tree_and_atomic_cross_wg>;
+
+    CGH.parallel_for<Name>(NDRange, Properties, [=](nd_item<Dims> NDIt) {
+      // Call user's functions. Reducer.MValue gets initialized there.
+      typename Reduction::reducer_type Reducer;
+      KernelFunc(NDIt, Reducer);
 
-    workGroupBarrier();
-    if (DoReducePartialSumsInLastWG[0]) {
-      // Reduce each result separately
-      // TODO: Opportunity to parallelize across elements
+      size_t WGSize = NDIt.get_local_range().size();
+      size_t LID = NDIt.get_local_linear_id();
+
+      // If there are multiple values, reduce each separately
+      // This prevents local memory from scaling with elements
       for (int E = 0; E < NElements; ++E) {
-        auto LocalSum = Identity;
-        for (size_t I = LID; I < NWorkGroups; I += WGSize)
-          LocalSum = BOp(LocalSum, PartialSums[I * NElements + E]);
 
-        LocalReds[LID] = LocalSum;
-        if (LID == 0)
-          LocalReds[WGSize] = Identity;
-        workGroupBarrier();
+        // Copy the element to local memory to prepare it for tree-reduction.
+        LocalReds[LID] = Reducer.getElement(E);
+        if (!IsPow2WG)
+          LocalReds[WGSize] = Reducer.getIdentity();
+        NDIt.barrier();
 
+        // Tree-reduction: reduce the local array LocalReds[:] to LocalReds[0].
+        // LocalReds[WGSize] accumulates last/odd elements when the step
+        // of tree-reduction loop is not even.
+        typename Reduction::binary_operation BOp;
         size_t PrevStep = WGSize;
         for (size_t CurStep = PrevStep >> 1; CurStep > 0; CurStep >>= 1) {
           if (LID < CurStep)
             LocalReds[LID] = BOp(LocalReds[LID], LocalReds[LID + CurStep]);
-          else if (LID == CurStep && (PrevStep & 0x1))
+          else if (!IsPow2WG && LID == CurStep && (PrevStep & 0x1))
             LocalReds[WGSize] = BOp(LocalReds[WGSize], LocalReds[PrevStep - 1]);
-          workGroupBarrier();
+          NDIt.barrier();
           PrevStep = CurStep;
         }
+
         if (LID == 0) {
-          auto V = BOp(LocalReds[0], LocalReds[WGSize]);
-          if (IsUpdateOfUserVar)
-            V = BOp(V, Out[E]);
-          Out[E] = V;
+          Reducer.getElement(E) =
+              IsPow2WG ? LocalReds[0] : BOp(LocalReds[0], LocalReds[WGSize]);
+        }
+
+        // Ensure item 0 is finished with LocalReds before next iteration
+        if (E != NElements - 1) {
+          NDIt.barrier();
         }
       }
-    }
-  });
 
-  if (Reduction::is_usm)
-    reduction::withAuxHandler(CGH, [&](handler &CopyHandler) {
-      reduSaveFinalResultToUserMem<KernelName>(CopyHandler, Redu);
+      if (LID == 0) {
+        Reducer.atomic_combine(&Out[0]);
+      }
     });
-}
 
-namespace reduction {
-namespace main_krn {
-template <class KernelName> struct NDRangeBothFastReduceAndAtomics;
-} // namespace main_krn
-} // namespace reduction
-/// Implements a command group function that enqueues a kernel that calls
-/// user's lambda function KernelFunc and also does one iteration of reduction
-/// of elements computed in user's lambda function.
-/// This version uses reduce() algorithm to reduce elements in each
-/// of work-groups, then it calls fast SYCL atomic operations to update
-/// the given reduction variable \p Out.
-///
-/// Briefly: calls user's lambda, reduce() + atomic, INT +
-/// ADD/MIN/MAX.
-template <typename KernelName, typename KernelType, int Dims,
-          typename PropertiesT, class Reduction>
-void reduCGFuncForNDRangeBothFastReduceAndAtomics(handler &CGH,
-                                                  KernelType KernelFunc,
-                                                  const nd_range<Dims> &Range,
-                                                  PropertiesT Properties,
-                                                  Reduction &Redu) {
-  auto Out = Redu.getReadWriteAccessorToInitializedMem(CGH);
-  size_t NElements = Reduction::num_elements;
-  using Name = __sycl_reduction_kernel<
-      reduction::main_krn::NDRangeBothFastReduceAndAtomics, KernelName>;
-  CGH.parallel_for<Name>(Range, Properties, [=](nd_item<Dims> NDIt) {
-    // Call user's function. Reducer.MValue gets initialized there.
-    typename Reduction::reducer_type Reducer;
-    KernelFunc(NDIt, Reducer);
-
-    typename Reduction::binary_operation BOp;
-    for (int E = 0; E < NElements; ++E) {
-      Reducer.getElement(E) =
-          reduce_over_group(NDIt.get_group(), Reducer.getElement(E), BOp);
+    if (Reduction::is_usm || Redu.initializeToIdentity()) {
+      reduction::withAuxHandler(CGH, [&](handler &CopyHandler) {
+        reduSaveFinalResultToUserMem<KernelName>(CopyHandler, Redu);
+      });
     }
-    if (NDIt.get_local_linear_id() == 0)
-      Reducer.atomic_combine(&Out[0]);
-  });
-
-  if (Reduction::is_usm || Redu.initializeToIdentity()) {
-    reduction::withAuxHandler(CGH, [&](handler &CopyHandler) {
-      reduSaveFinalResultToUserMem<KernelName>(CopyHandler, Redu);
-    });
   }
-}
+};
 
-namespace reduction {
-namespace main_krn {
-template <class KernelName> struct NDRangeFastAtomicsOnly;
-} // namespace main_krn
-} // namespace reduction
-/// Implements a command group function that enqueues a kernel that calls
-/// user's lambda function KernelFunc and also does one iteration of reduction
-/// of elements computed in user's lambda function.
-/// This version uses tree-reduction algorithm to reduce elements in each
-/// of work-groups, then it calls fast SYCL atomic operations to update
-/// user's reduction variable.
-///
-/// Briefly: calls user's lambda, tree-reduction + atomic, INT + AND/OR/XOR.
-template <typename KernelName, typename KernelType, int Dims,
-          typename PropertiesT, class Reduction>
-void reduCGFuncForNDRangeFastAtomicsOnly(handler &CGH, KernelType KernelFunc,
-                                         const nd_range<Dims> &Range,
-                                         PropertiesT Properties,
-                                         Reduction &Redu) {
-  auto Out = Redu.getReadWriteAccessorToInitializedMem(CGH);
-  size_t NElements = Reduction::num_elements;
-  size_t WGSize = Range.get_local_range().size();
-  bool IsPow2WG = (WGSize & (WGSize - 1)) == 0;
-
-  // Use local memory to reduce elements in work-groups into zero-th element.
-  // If WGSize is not power of two, then WGSize+1 elements are allocated.
-  // The additional last element is used to catch reduce elements that could
-  // otherwise be lost in the tree-reduction algorithm used in the kernel.
-  size_t NLocalElements = WGSize + (IsPow2WG ? 0 : 1);
-  local_accessor<typename Reduction::result_type, 1> LocalReds{NLocalElements,
-                                                               CGH};
-
-  using Name =
-      __sycl_reduction_kernel<reduction::main_krn::NDRangeFastAtomicsOnly,
-                              KernelName>;
-  CGH.parallel_for<Name>(Range, Properties, [=](nd_item<Dims> NDIt) {
-    // Call user's functions. Reducer.MValue gets initialized there.
-    typename Reduction::reducer_type Reducer;
-    KernelFunc(NDIt, Reducer);
-
-    size_t WGSize = NDIt.get_local_range().size();
-    size_t LID = NDIt.get_local_linear_id();
-
-    // If there are multiple values, reduce each separately
-    // This prevents local memory from scaling with elements
-    for (int E = 0; E < NElements; ++E) {
-
-      // Copy the element to local memory to prepare it for tree-reduction.
-      LocalReds[LID] = Reducer.getElement(E);
-      if (!IsPow2WG)
-        LocalReds[WGSize] = Reducer.getIdentity();
-      NDIt.barrier();
+template <>
+struct NDRangeReduction<
+    reduction::strategy::group_reduce_and_multiple_kernels> {
+  template <typename KernelName, int Dims, typename PropertiesT,
+            typename KernelType, typename Reduction>
+  static void run(handler &CGH, std::shared_ptr<detail::queue_impl> &Queue,
+                  nd_range<Dims> NDRange, PropertiesT &Properties,
+                  Reduction &Redu, KernelType &KernelFunc) {
+    // Before running the kernels, check that device has enough local memory
+    // to hold local arrays that may be required for the reduction algorithm.
+    // TODO: If the work-group-size is limited by the local memory, then
+    // a special version of the main kernel may be created. The one that would
+    // not use local accessors, which means it would not do the reduction in
+    // the main kernel, but simply generate Range.get_global_range.size() number
+    // of partial sums, leaving the reduction work to the additional/aux
+    // kernels.
+    constexpr bool HFR = Reduction::has_fast_reduce;
+    size_t OneElemSize = HFR ? 0 : sizeof(typename Reduction::result_type);
+    // TODO: currently the maximal work group size is determined for the given
+    // queue/device, while it may be safer to use queries to the kernel compiled
+    // for the device.
+    size_t MaxWGSize = reduGetMaxWGSize(Queue, OneElemSize);
+    if (NDRange.get_local_range().size() > MaxWGSize)
+      throw sycl::runtime_error("The implementation handling parallel_for with"
+                                " reduction requires work group size not bigger"
+                                " than " +
+                                    std::to_string(MaxWGSize),
+                                PI_ERROR_INVALID_WORK_GROUP_SIZE);
+
+    size_t NElements = Reduction::num_elements;
+    size_t NWorkGroups = NDRange.get_group_range().size();
+    auto Out = Redu.getWriteAccForPartialReds(NWorkGroups * NElements, CGH);
+
+    bool IsUpdateOfUserVar =
+        !Reduction::is_usm && !Redu.initializeToIdentity() && NWorkGroups == 1;
+
+    using Name = __sycl_reduction_kernel<
+        reduction::MainKrn, KernelName,
+        reduction::strategy::group_reduce_and_multiple_kernels>;
+
+    CGH.parallel_for<Name>(NDRange, Properties, [=](nd_item<Dims> NDIt) {
+      // Call user's functions. Reducer.MValue gets initialized there.
+      typename Reduction::reducer_type Reducer;
+      KernelFunc(NDIt, Reducer);
 
-      // Tree-reduction: reduce the local array LocalReds[:] to LocalReds[0].
-      // LocalReds[WGSize] accumulates last/odd elements when the step
-      // of tree-reduction loop is not even.
+      // Compute the partial sum/reduction for the work-group.
+      size_t WGID = NDIt.get_group_linear_id();
       typename Reduction::binary_operation BOp;
-      size_t PrevStep = WGSize;
-      for (size_t CurStep = PrevStep >> 1; CurStep > 0; CurStep >>= 1) {
-        if (LID < CurStep)
-          LocalReds[LID] = BOp(LocalReds[LID], LocalReds[LID + CurStep]);
-        else if (!IsPow2WG && LID == CurStep && (PrevStep & 0x1))
-          LocalReds[WGSize] = BOp(LocalReds[WGSize], LocalReds[PrevStep - 1]);
-        NDIt.barrier();
-        PrevStep = CurStep;
-      }
-
-      if (LID == 0) {
-        Reducer.getElement(E) =
-            IsPow2WG ? LocalReds[0] : BOp(LocalReds[0], LocalReds[WGSize]);
-      }
-
-      // Ensure item 0 is finished with LocalReds before next iteration
-      if (E != NElements - 1) {
-        NDIt.barrier();
+      for (int E = 0; E < NElements; ++E) {
+        typename Reduction::result_type PSum;
+        PSum = Reducer.getElement(E);
+        PSum = reduce_over_group(NDIt.get_group(), PSum, BOp);
+        if (NDIt.get_local_linear_id() == 0) {
+          if (IsUpdateOfUserVar)
+            PSum = BOp(Out[E], PSum);
+          Out[WGID * NElements + E] = PSum;
+        }
       }
-    }
-
-    if (LID == 0) {
-      Reducer.atomic_combine(&Out[0]);
-    }
-  });
-
-  if (Reduction::is_usm || Redu.initializeToIdentity()) {
-    reduction::withAuxHandler(CGH, [&](handler &CopyHandler) {
-      reduSaveFinalResultToUserMem<KernelName>(CopyHandler, Redu);
     });
-  }
-}
 
-namespace reduction {
-namespace main_krn {
-template <class KernelName> struct NDRangeFastReduceOnly;
-} // namespace main_krn
-} // namespace reduction
-/// Implements a command group function that enqueues a kernel that
-/// calls user's lambda function and does one iteration of reduction
-/// of elements in each of work-groups.
-/// This version uses reduce() algorithm to reduce elements in each
-/// of work-groups. At the end of each work-groups the partial sum is written
-/// to a global buffer.
-///
-/// Briefly: user's lambda, reduce(), FP + ADD/MIN/MAX.
-template <typename KernelName, typename KernelType, int Dims,
-          typename PropertiesT, class Reduction>
-void reduCGFuncForNDRangeFastReduceOnly(handler &CGH, KernelType KernelFunc,
-                                        const nd_range<Dims> &Range,
-                                        PropertiesT Properties,
-                                        Reduction &Redu) {
-  size_t NElements = Reduction::num_elements;
-  size_t NWorkGroups = Range.get_group_range().size();
-  auto Out = Redu.getWriteAccForPartialReds(NWorkGroups * NElements, CGH);
-
-  bool IsUpdateOfUserVar =
-      !Reduction::is_usm && !Redu.initializeToIdentity() && NWorkGroups == 1;
-
-  using Name =
-      __sycl_reduction_kernel<reduction::main_krn::NDRangeFastReduceOnly,
-                              KernelName>;
-  CGH.parallel_for<Name>(Range, Properties, [=](nd_item<Dims> NDIt) {
-    // Call user's functions. Reducer.MValue gets initialized there.
-    typename Reduction::reducer_type Reducer;
-    KernelFunc(NDIt, Reducer);
-
-    // Compute the partial sum/reduction for the work-group.
-    size_t WGID = NDIt.get_group_linear_id();
-    typename Reduction::binary_operation BOp;
-    for (int E = 0; E < NElements; ++E) {
-      typename Reduction::result_type PSum;
-      PSum = Reducer.getElement(E);
-      PSum = reduce_over_group(NDIt.get_group(), PSum, BOp);
-      if (NDIt.get_local_linear_id() == 0) {
-        if (IsUpdateOfUserVar)
-          PSum = BOp(Out[E], PSum);
-        Out[WGID * NElements + E] = PSum;
-      }
+    reduction::finalizeHandler(CGH);
+
+    // Run the additional kernel as many times as needed to reduce all partial
+    // sums into one scalar.
+
+    // TODO: Create a special slow/sequential version of the kernel that would
+    // handle the reduction instead of reporting an assert below.
+    if (MaxWGSize <= 1)
+      throw sycl::runtime_error("The implementation handling parallel_for with "
+                                "reduction requires the maximal work group "
+                                "size to be greater than 1 to converge. "
+                                "The maximal work group size depends on the "
+                                "device and the size of the objects passed to "
+                                "the reduction.",
+                                PI_ERROR_INVALID_WORK_GROUP_SIZE);
+    size_t NWorkItems = NDRange.get_group_range().size();
+    while (NWorkItems > 1) {
+      reduction::withAuxHandler(CGH, [&](handler &AuxHandler) {
+        size_t NElements = Reduction::num_elements;
+        size_t NWorkGroups;
+        size_t WGSize = reduComputeWGSize(NWorkItems, MaxWGSize, NWorkGroups);
+
+        // The last work-group may be not fully loaded with work, or the work
+        // group size may be not power of two. Those two cases considered
+        // inefficient as they require additional code and checks in the kernel.
+        bool HasUniformWG = NWorkGroups * WGSize == NWorkItems;
+        if (!Reduction::has_fast_reduce)
+          HasUniformWG = HasUniformWG && (WGSize & (WGSize - 1)) == 0;
+
+        // Get read accessor to the buffer that was used as output
+        // in the previous kernel.
+        auto In = Redu.getReadAccToPreviousPartialReds(AuxHandler);
+        auto Out =
+            Redu.getWriteAccForPartialReds(NWorkGroups * NElements, AuxHandler);
+
+        using Name = __sycl_reduction_kernel<
+            reduction::AuxKrn, KernelName,
+            reduction::strategy::group_reduce_and_multiple_kernels>;
+
+        bool IsUpdateOfUserVar = !Reduction::is_usm &&
+                                 !Redu.initializeToIdentity() &&
+                                 NWorkGroups == 1;
+        range<1> GlobalRange = {HasUniformWG ? NWorkItems
+                                             : NWorkGroups * WGSize};
+        nd_range<1> Range{GlobalRange, range<1>(WGSize)};
+        AuxHandler.parallel_for<Name>(Range, [=](nd_item<1> NDIt) {
+          typename Reduction::binary_operation BOp;
+          size_t WGID = NDIt.get_group_linear_id();
+          size_t GID = NDIt.get_global_linear_id();
+
+          for (int E = 0; E < NElements; ++E) {
+            typename Reduction::result_type PSum =
+                (HasUniformWG || (GID < NWorkItems))
+                    ? In[GID * NElements + E]
+                    : Reduction::reducer_type::getIdentity();
+            PSum = reduce_over_group(NDIt.get_group(), PSum, BOp);
+            if (NDIt.get_local_linear_id() == 0) {
+              if (IsUpdateOfUserVar)
+                PSum = BOp(Out[E], PSum);
+              Out[WGID * NElements + E] = PSum;
+            }
+          }
+        });
+        NWorkItems = NWorkGroups;
+      });
+    } // end while (NWorkItems > 1)
+
+    if (Reduction::is_usm) {
+      reduction::withAuxHandler(CGH, [&](handler &CopyHandler) {
+        reduSaveFinalResultToUserMem<KernelName>(CopyHandler, Redu);
+      });
     }
-  });
-}
+  }
+};
 
-namespace reduction {
-namespace main_krn {
-template <class KernelName> struct NDRangeBasic;
-} // namespace main_krn
-} // namespace reduction
-/// Implements a command group function that enqueues a kernel that calls
-/// user's lambda function \param KernelFunc and does one iteration of
-/// reduction of elements in each of work-groups.
-/// This version uses tree-reduction algorithm to reduce elements in each
-/// of work-groups. At the end of each work-group the partial sum is written
-/// to a global buffer.
-///
-/// Briefly: user's lambda, tree-reduction, CUSTOM types/ops.
-template <typename KernelName, typename KernelType, int Dims,
-          typename PropertiesT, class Reduction>
-void reduCGFuncForNDRangeBasic(handler &CGH, KernelType KernelFunc,
-                               const nd_range<Dims> &Range,
-                               PropertiesT Properties, Reduction &Redu) {
-  size_t NElements = Reduction::num_elements;
-  size_t WGSize = Range.get_local_range().size();
-  bool IsPow2WG = (WGSize & (WGSize - 1)) == 0;
-  size_t NWorkGroups = Range.get_group_range().size();
-  auto Out = Redu.getWriteAccForPartialReds(NWorkGroups * NElements, CGH);
-
-  bool IsUpdateOfUserVar =
-      !Reduction::is_usm && !Redu.initializeToIdentity() && NWorkGroups == 1;
-
-  // Use local memory to reduce elements in work-groups into 0-th element.
-  // If WGSize is not power of two, then WGSize+1 elements are allocated.
-  // The additional last element is used to catch elements that could
-  // otherwise be lost in the tree-reduction algorithm.
-  size_t NumLocalElements = WGSize + (IsPow2WG ? 0 : 1);
-  local_accessor<typename Reduction::result_type, 1> LocalReds{NumLocalElements,
-                                                               CGH};
-  typename Reduction::result_type ReduIdentity = Redu.getIdentity();
-  using Name =
-      __sycl_reduction_kernel<reduction::main_krn::NDRangeBasic, KernelName>;
-  auto BOp = Redu.getBinaryOperation();
-  CGH.parallel_for<Name>(Range, Properties, [=](nd_item<Dims> NDIt) {
-    // Call user's functions. Reducer.MValue gets initialized there.
-    typename Reduction::reducer_type Reducer(ReduIdentity, BOp);
-    KernelFunc(NDIt, Reducer);
-
-    size_t WGSize = NDIt.get_local_range().size();
-    size_t LID = NDIt.get_local_linear_id();
-
-    // If there are multiple values, reduce each separately
-    // This prevents local memory from scaling with elements
-    for (int E = 0; E < NElements; ++E) {
-
-      // Copy the element to local memory to prepare it for tree-reduction.
-      LocalReds[LID] = Reducer.getElement(E);
-      if (!IsPow2WG)
-        LocalReds[WGSize] = ReduIdentity;
-      NDIt.barrier();
+template <> struct NDRangeReduction<reduction::strategy::basic> {
+  template <typename KernelName, int Dims, typename PropertiesT,
+            typename KernelType, typename Reduction>
+  static void run(handler &CGH, std::shared_ptr<detail::queue_impl> &Queue,
+                  nd_range<Dims> NDRange, PropertiesT &Properties,
+                  Reduction &Redu, KernelType &KernelFunc) {
+    constexpr bool HFR = Reduction::has_fast_reduce;
+    size_t OneElemSize = HFR ? 0 : sizeof(typename Reduction::result_type);
+    // TODO: currently the maximal work group size is determined for the given
+    // queue/device, while it may be safer to use queries to the kernel compiled
+    // for the device.
+    size_t MaxWGSize = reduGetMaxWGSize(Queue, OneElemSize);
+    if (NDRange.get_local_range().size() > MaxWGSize)
+      throw sycl::runtime_error("The implementation handling parallel_for with"
+                                " reduction requires work group size not bigger"
+                                " than " +
+                                    std::to_string(MaxWGSize),
+                                PI_ERROR_INVALID_WORK_GROUP_SIZE);
+
+    size_t NElements = Reduction::num_elements;
+    size_t WGSize = NDRange.get_local_range().size();
+    bool IsPow2WG = (WGSize & (WGSize - 1)) == 0;
+    size_t NWorkGroups = NDRange.get_group_range().size();
+    auto Out = Redu.getWriteAccForPartialReds(NWorkGroups * NElements, CGH);
+
+    bool IsUpdateOfUserVar =
+        !Reduction::is_usm && !Redu.initializeToIdentity() && NWorkGroups == 1;
+
+    // Use local memory to reduce elements in work-groups into 0-th element.
+    // If WGSize is not power of two, then WGSize+1 elements are allocated.
+    // The additional last element is used to catch elements that could
+    // otherwise be lost in the tree-reduction algorithm.
+    size_t NumLocalElements = WGSize + (IsPow2WG ? 0 : 1);
+    local_accessor<typename Reduction::result_type, 1> LocalReds{
+        NumLocalElements, CGH};
+    typename Reduction::result_type ReduIdentity = Redu.getIdentity();
+    using Name = __sycl_reduction_kernel<reduction::MainKrn, KernelName,
+                                         reduction::strategy::basic>;
+
+    auto BOp = Redu.getBinaryOperation();
+    CGH.parallel_for<Name>(NDRange, Properties, [=](nd_item<Dims> NDIt) {
+      // Call user's functions. Reducer.MValue gets initialized there.
+      typename Reduction::reducer_type Reducer(ReduIdentity, BOp);
+      KernelFunc(NDIt, Reducer);
 
-      // Tree-reduction: reduce the local array LocalReds[:] to LocalReds[0]
-      // LocalReds[WGSize] accumulates last/odd elements when the step
-      // of tree-reduction loop is not even.
-      size_t PrevStep = WGSize;
-      for (size_t CurStep = PrevStep >> 1; CurStep > 0; CurStep >>= 1) {
-        if (LID < CurStep)
-          LocalReds[LID] = BOp(LocalReds[LID], LocalReds[LID + CurStep]);
-        else if (!IsPow2WG && LID == CurStep && (PrevStep & 0x1))
-          LocalReds[WGSize] = BOp(LocalReds[WGSize], LocalReds[PrevStep - 1]);
-        NDIt.barrier();
-        PrevStep = CurStep;
-      }
+      size_t WGSize = NDIt.get_local_range().size();
+      size_t LID = NDIt.get_local_linear_id();
 
-      // Compute the partial sum/reduction for the work-group.
-      if (LID == 0) {
-        size_t GrID = NDIt.get_group_linear_id();
-        typename Reduction::result_type PSum =
-            IsPow2WG ? LocalReds[0] : BOp(LocalReds[0], LocalReds[WGSize]);
-        if (IsUpdateOfUserVar)
-          PSum = BOp(Out[0], PSum);
-        Out[GrID * NElements + E] = PSum;
-      }
+      // If there are multiple values, reduce each separately
+      // This prevents local memory from scaling with elements
+      for (int E = 0; E < NElements; ++E) {
 
-      // Ensure item 0 is finished with LocalReds before next iteration
-      if (E != NElements - 1) {
+        // Copy the element to local memory to prepare it for tree-reduction.
+        LocalReds[LID] = Reducer.getElement(E);
+        if (!IsPow2WG)
+          LocalReds[WGSize] = ReduIdentity;
         NDIt.barrier();
-      }
-    }
-  });
-}
 
-namespace reduction {
-namespace aux_krn {
-template <class KernelName> struct FastReduce;
-} // namespace aux_krn
-} // namespace reduction
-/// Implements a command group function that enqueues a kernel that does one
-/// iteration of reduction of elements in each of work-groups.
-/// This version uses reduce() algorithm to reduce elements in each
-/// of work-groups. At the end of each work-groups the partial sum is written
-/// to a global buffer.
-///
-/// Briefly: aux kernel, reduce(), reproducible results, FP +
-/// ADD/MIN/MAX
-template <typename KernelName, typename KernelType, class Reduction,
-          typename InputT, typename OutputT>
-void reduAuxCGFuncFastReduceImpl(handler &CGH, bool UniformWG,
-                                 size_t NWorkItems, size_t NWorkGroups,
-                                 size_t WGSize, Reduction &Redu, InputT In,
-                                 OutputT Out) {
-  size_t NElements = Reduction::num_elements;
-  using Name =
-      __sycl_reduction_kernel<reduction::aux_krn::FastReduce, KernelName>;
-  bool IsUpdateOfUserVar =
-      !Reduction::is_usm && !Redu.initializeToIdentity() && NWorkGroups == 1;
-  range<1> GlobalRange = {UniformWG ? NWorkItems : NWorkGroups * WGSize};
-  nd_range<1> Range{GlobalRange, range<1>(WGSize)};
-  CGH.parallel_for<Name>(Range, [=](nd_item<1> NDIt) {
-    typename Reduction::binary_operation BOp;
-    size_t WGID = NDIt.get_group_linear_id();
-    size_t GID = NDIt.get_global_linear_id();
-
-    for (int E = 0; E < NElements; ++E) {
-      typename Reduction::result_type PSum =
-          (UniformWG || (GID < NWorkItems))
-              ? In[GID * NElements + E]
-              : Reduction::reducer_type::getIdentity();
-      PSum = reduce_over_group(NDIt.get_group(), PSum, BOp);
-      if (NDIt.get_local_linear_id() == 0) {
-        if (IsUpdateOfUserVar)
-          PSum = BOp(Out[E], PSum);
-        Out[WGID * NElements + E] = PSum;
-      }
-    }
-  });
-}
+        // Tree-reduction: reduce the local array LocalReds[:] to LocalReds[0]
+        // LocalReds[WGSize] accumulates last/odd elements when the step
+        // of tree-reduction loop is not even.
+        size_t PrevStep = WGSize;
+        for (size_t CurStep = PrevStep >> 1; CurStep > 0; CurStep >>= 1) {
+          if (LID < CurStep)
+            LocalReds[LID] = BOp(LocalReds[LID], LocalReds[LID + CurStep]);
+          else if (!IsPow2WG && LID == CurStep && (PrevStep & 0x1))
+            LocalReds[WGSize] = BOp(LocalReds[WGSize], LocalReds[PrevStep - 1]);
+          NDIt.barrier();
+          PrevStep = CurStep;
+        }
 
-namespace reduction {
-namespace aux_krn {
-template <class KernelName> struct NoFastReduceNorAtomic;
-} // namespace aux_krn
-} // namespace reduction
-/// Implements a command group function that enqueues a kernel that does one
-/// iteration of reduction of elements in each of work-groups.
-/// This version uses tree-reduction algorithm to reduce elements in each
-/// of work-groups. At the end of each work-group the partial sum is written
-/// to a global buffer.
-///
-/// Briefly: aux kernel, tree-reduction, CUSTOM types/ops.
-template <typename KernelName, typename KernelType, class Reduction,
-          typename InputT, typename OutputT>
-void reduAuxCGFuncNoFastReduceNorAtomicImpl(handler &CGH, bool UniformPow2WG,
-                                            size_t NWorkItems,
-                                            size_t NWorkGroups, size_t WGSize,
-                                            Reduction &Redu, InputT In,
-                                            OutputT Out) {
-  size_t NElements = Reduction::num_elements;
-  bool IsUpdateOfUserVar =
-      !Reduction::is_usm && !Redu.initializeToIdentity() && NWorkGroups == 1;
-
-  // Use local memory to reduce elements in work-groups into 0-th element.
-  // If WGSize is not power of two, then WGSize+1 elements are allocated.
-  // The additional last element is used to catch elements that could
-  // otherwise be lost in the tree-reduction algorithm.
-  size_t NumLocalElements = WGSize + (UniformPow2WG ? 0 : 1);
-  local_accessor<typename Reduction::result_type, 1> LocalReds{NumLocalElements,
-                                                               CGH};
-
-  auto ReduIdentity = Redu.getIdentity();
-  auto BOp = Redu.getBinaryOperation();
-  using Name =
-      __sycl_reduction_kernel<reduction::aux_krn::NoFastReduceNorAtomic,
-                              KernelName>;
-  range<1> GlobalRange = {UniformPow2WG ? NWorkItems : NWorkGroups * WGSize};
-  nd_range<1> Range{GlobalRange, range<1>(WGSize)};
-  CGH.parallel_for<Name>(Range, [=](nd_item<1> NDIt) {
-    size_t WGSize = NDIt.get_local_range().size();
-    size_t LID = NDIt.get_local_linear_id();
-    size_t GID = NDIt.get_global_linear_id();
-
-    for (int E = 0; E < NElements; ++E) {
-      // Copy the element to local memory to prepare it for tree-reduction.
-      LocalReds[LID] = (UniformPow2WG || GID < NWorkItems)
-                           ? In[GID * NElements + E]
-                           : ReduIdentity;
-      if (!UniformPow2WG)
-        LocalReds[WGSize] = ReduIdentity;
-      NDIt.barrier();
+        // Compute the partial sum/reduction for the work-group.
+        if (LID == 0) {
+          size_t GrID = NDIt.get_group_linear_id();
+          typename Reduction::result_type PSum =
+              IsPow2WG ? LocalReds[0] : BOp(LocalReds[0], LocalReds[WGSize]);
+          if (IsUpdateOfUserVar)
+            PSum = BOp(Out[0], PSum);
+          Out[GrID * NElements + E] = PSum;
+        }
 
-      // Tree-reduction: reduce the local array LocalReds[:] to LocalReds[0]
-      // LocalReds[WGSize] accumulates last/odd elements when the step
-      // of tree-reduction loop is not even.
-      size_t PrevStep = WGSize;
-      for (size_t CurStep = PrevStep >> 1; CurStep > 0; CurStep >>= 1) {
-        if (LID < CurStep)
-          LocalReds[LID] = BOp(LocalReds[LID], LocalReds[LID + CurStep]);
-        else if (!UniformPow2WG && LID == CurStep && (PrevStep & 0x1))
-          LocalReds[WGSize] = BOp(LocalReds[WGSize], LocalReds[PrevStep - 1]);
-        NDIt.barrier();
-        PrevStep = CurStep;
+        // Ensure item 0 is finished with LocalReds before next iteration
+        if (E != NElements - 1) {
+          NDIt.barrier();
+        }
       }
+    });
 
-      // Compute the partial sum/reduction for the work-group.
-      if (LID == 0) {
-        size_t GrID = NDIt.get_group_linear_id();
-        typename Reduction::result_type PSum =
-            UniformPow2WG ? LocalReds[0] : BOp(LocalReds[0], LocalReds[WGSize]);
-        if (IsUpdateOfUserVar)
-          PSum = BOp(Out[0], PSum);
-        Out[GrID * NElements + E] = PSum;
-      }
+    reduction::finalizeHandler(CGH);
+
+    // 2. Run the additional kernel as many times as needed to reduce
+    // all partial sums into one scalar.
+
+    // TODO: Create a special slow/sequential version of the kernel that would
+    // handle the reduction instead of reporting an assert below.
+    if (MaxWGSize <= 1)
+      throw sycl::runtime_error("The implementation handling parallel_for with "
+                                "reduction requires the maximal work group "
+                                "size to be greater than 1 to converge. "
+                                "The maximal work group size depends on the "
+                                "device and the size of the objects passed to "
+                                "the reduction.",
+                                PI_ERROR_INVALID_WORK_GROUP_SIZE);
+    size_t NWorkItems = NDRange.get_group_range().size();
+    while (NWorkItems > 1) {
+      reduction::withAuxHandler(CGH, [&](handler &AuxHandler) {
+        size_t NElements = Reduction::num_elements;
+        size_t NWorkGroups;
+        size_t WGSize = reduComputeWGSize(NWorkItems, MaxWGSize, NWorkGroups);
+
+        // The last work-group may be not fully loaded with work, or the work
+        // group size may be not power of two. Those two cases considered
+        // inefficient as they require additional code and checks in the kernel.
+        bool HasUniformWG = NWorkGroups * WGSize == NWorkItems;
+        if (!Reduction::has_fast_reduce)
+          HasUniformWG = HasUniformWG && (WGSize & (WGSize - 1)) == 0;
+
+        // Get read accessor to the buffer that was used as output
+        // in the previous kernel.
+        auto In = Redu.getReadAccToPreviousPartialReds(AuxHandler);
+        auto Out =
+            Redu.getWriteAccForPartialReds(NWorkGroups * NElements, AuxHandler);
+
+        bool IsUpdateOfUserVar = !Reduction::is_usm &&
+                                 !Redu.initializeToIdentity() &&
+                                 NWorkGroups == 1;
+
+        bool UniformPow2WG = HasUniformWG;
+        // Use local memory to reduce elements in work-groups into 0-th element.
+        // If WGSize is not power of two, then WGSize+1 elements are allocated.
+        // The additional last element is used to catch elements that could
+        // otherwise be lost in the tree-reduction algorithm.
+        size_t NumLocalElements = WGSize + (UniformPow2WG ? 0 : 1);
+        local_accessor<typename Reduction::result_type, 1> LocalReds{
+            NumLocalElements, AuxHandler};
+
+        auto ReduIdentity = Redu.getIdentity();
+        auto BOp = Redu.getBinaryOperation();
+        using Name = __sycl_reduction_kernel<reduction::AuxKrn, KernelName,
+                                             reduction::strategy::basic>;
+
+        range<1> GlobalRange = {UniformPow2WG ? NWorkItems
+                                              : NWorkGroups * WGSize};
+        nd_range<1> Range{GlobalRange, range<1>(WGSize)};
+        AuxHandler.parallel_for<Name>(Range, [=](nd_item<1> NDIt) {
+          size_t WGSize = NDIt.get_local_range().size();
+          size_t LID = NDIt.get_local_linear_id();
+          size_t GID = NDIt.get_global_linear_id();
+
+          for (int E = 0; E < NElements; ++E) {
+            // Copy the element to local memory to prepare it for
+            // tree-reduction.
+            LocalReds[LID] = (UniformPow2WG || GID < NWorkItems)
+                                 ? In[GID * NElements + E]
+                                 : ReduIdentity;
+            if (!UniformPow2WG)
+              LocalReds[WGSize] = ReduIdentity;
+            NDIt.barrier();
+
+            // Tree-reduction: reduce the local array LocalReds[:] to
+            // LocalReds[0] LocalReds[WGSize] accumulates last/odd elements when
+            // the step of tree-reduction loop is not even.
+            size_t PrevStep = WGSize;
+            for (size_t CurStep = PrevStep >> 1; CurStep > 0; CurStep >>= 1) {
+              if (LID < CurStep)
+                LocalReds[LID] = BOp(LocalReds[LID], LocalReds[LID + CurStep]);
+              else if (!UniformPow2WG && LID == CurStep && (PrevStep & 0x1))
+                LocalReds[WGSize] =
+                    BOp(LocalReds[WGSize], LocalReds[PrevStep - 1]);
+              NDIt.barrier();
+              PrevStep = CurStep;
+            }
+
+            // Compute the partial sum/reduction for the work-group.
+            if (LID == 0) {
+              size_t GrID = NDIt.get_group_linear_id();
+              typename Reduction::result_type PSum =
+                  UniformPow2WG ? LocalReds[0]
+                                : BOp(LocalReds[0], LocalReds[WGSize]);
+              if (IsUpdateOfUserVar)
+                PSum = BOp(Out[0], PSum);
+              Out[GrID * NElements + E] = PSum;
+            }
+
+            // Ensure item 0 is finished with LocalReds before next iteration
+            if (E != NElements - 1) {
+              NDIt.barrier();
+            }
+          }
+        });
+        NWorkItems = NWorkGroups;
+      });
+    } // end while (NWorkItems > 1)
+
+    if (Reduction::is_usm) {
+      reduction::withAuxHandler(CGH, [&](handler &CopyHandler) {
+        reduSaveFinalResultToUserMem<KernelName>(CopyHandler, Redu);
+      });
+    }
+  }
+};
 
-      // Ensure item 0 is finished with LocalReds before next iteration
-      if (E != NElements - 1) {
-        NDIt.barrier();
+// Auto-dispatch. Must be the last one.
+template <> struct NDRangeReduction<reduction::strategy::auto_select> {
+  // Some readability aliases, to increase signal/noise ratio below.
+  template <reduction::strategy Strategy>
+  using Impl = NDRangeReduction<Strategy>;
+  using S = reduction::strategy;
+
+  template <typename KernelName, int Dims, typename PropertiesT,
+            typename KernelType, typename Reduction>
+  static void run(handler &CGH, std::shared_ptr<detail::queue_impl> &Queue,
+                  nd_range<Dims> NDRange, PropertiesT &Properties,
+                  Reduction &Redu, KernelType &KernelFunc) {
+    auto Delegate = [&](auto Impl) {
+      Impl.template run<KernelName>(CGH, Queue, NDRange, Properties, Redu,
+                                    KernelFunc);
+    };
+
+    if constexpr (Reduction::has_float64_atomics) {
+      if (getDeviceFromHandler(CGH).has(aspect::atomic64))
+        return Delegate(Impl<S::group_reduce_and_atomic_cross_wg>{});
+
+      if constexpr (Reduction::has_fast_reduce)
+        return Delegate(Impl<S::group_reduce_and_multiple_kernels>{});
+      else
+        return Delegate(Impl<S::basic>{});
+    } else if constexpr (Reduction::has_fast_atomics) {
+      if constexpr (Reduction::has_fast_reduce) {
+        return Delegate(Impl<S::group_reduce_and_atomic_cross_wg>{});
+      } else {
+        return Delegate(Impl<S::local_mem_tree_and_atomic_cross_wg>{});
       }
+    } else {
+      if constexpr (Reduction::has_fast_reduce)
+        return Delegate(Impl<S::group_reduce_and_multiple_kernels>{});
+      else
+        return Delegate(Impl<S::basic>{});
     }
-  });
-}
-
-/// Implements a command group function that enqueues a kernel that does one
-/// iteration of reduction of elements in each of work-groups.
-/// At the end of each work-group the partial sum is written to a global buffer.
-/// The function returns the number of the newly generated partial sums.
-template <typename KernelName, typename KernelType, class Reduction>
-size_t reduAuxCGFunc(handler &CGH, size_t NWorkItems, size_t MaxWGSize,
-                     Reduction &Redu) {
-  constexpr size_t NElements = Reduction::num_elements;
-  size_t NWorkGroups;
-  size_t WGSize = reduComputeWGSize(NWorkItems, MaxWGSize, NWorkGroups);
-
-  // The last work-group may be not fully loaded with work, or the work group
-  // size may be not power of two. Those two cases considered inefficient
-  // as they require additional code and checks in the kernel.
-  bool HasUniformWG = NWorkGroups * WGSize == NWorkItems;
-  if (!Reduction::has_fast_reduce)
-    HasUniformWG = HasUniformWG && (WGSize & (WGSize - 1)) == 0;
-
-  // Get read accessor to the buffer that was used as output
-  // in the previous kernel.
-  auto In = Redu.getReadAccToPreviousPartialReds(CGH);
-  auto Out = Redu.getWriteAccForPartialReds(NWorkGroups * NElements, CGH);
-
-  if constexpr (Reduction::has_fast_reduce && !Reduction::has_fast_atomics) {
-    reduAuxCGFuncFastReduceImpl<KernelName, KernelType>(
-        CGH, HasUniformWG, NWorkItems, NWorkGroups, WGSize, Redu, In, Out);
 
-  } else {
-    reduAuxCGFuncNoFastReduceNorAtomicImpl<KernelName, KernelType>(
-        CGH, HasUniformWG, NWorkItems, NWorkGroups, WGSize, Redu, In, Out);
+    assert(false && "Must be unreachable!");
   }
-  return NWorkGroups;
-}
+};
 
 /// For the given 'Reductions' types pack and indices enumerating them this
 /// function either creates new temporary accessors for partial sums (if IsOneWG
@@ -1852,8 +1926,10 @@ void reduCGFuncMulti(handler &CGH, KernelType KernelFunc,
     std::array InitToIdentityProps{
         std::get<Is>(ReduTuple).initializeToIdentity()...};
 
-    using Name = __sycl_reduction_kernel<reduction::main_krn::NDRangeMulti,
-                                         KernelName, decltype(OutAccsTuple)>;
+    using Name = __sycl_reduction_kernel<reduction::MainKrn, KernelName,
+                                         reduction::strategy::multi,
+                                         decltype(OutAccsTuple)>;
+
     CGH.parallel_for<Name>(Range, Properties, [=](nd_item<Dims> NDIt) {
       // Pass all reductions to user's lambda in the same order as supplied
       // Each reducer initializes its own storage
@@ -2095,7 +2171,8 @@ size_t reduAuxCGFunc(handler &CGH, size_t NWorkItems, size_t MaxWGSize,
   auto Rest = [&](auto Predicate, auto OutAccsTuple) {
     auto AccReduIndices = filterSequence<Reductions...>(Predicate, ReduIndices);
     associateReduAccsWithHandler(CGH, ReduTuple, AccReduIndices);
-    using Name = __sycl_reduction_kernel<reduction::aux_krn::Multi, KernelName,
+    using Name = __sycl_reduction_kernel<reduction::AuxKrn, KernelName,
+                                         reduction::strategy::multi,
                                          decltype(Predicate)>;
     // TODO: Opportunity to parallelize across number of elements
     range<1> GlobalRange = {HasUniformWG ? NWorkItems : NWorkGroups * WGSize};
@@ -2150,6 +2227,16 @@ tuple_select_elements(TupleT Tuple, std::index_sequence<Is...>) {
   return {std::get<Is>(std::move(Tuple))...};
 }
 
+template <typename KernelName, reduction::strategy Strategy, int Dims,
+          typename PropertiesT, typename KernelType, typename Reduction>
+void reduction_parallel_for(handler &CGH,
+                            std::shared_ptr<detail::queue_impl> Queue,
+                            nd_range<Dims> NDRange, PropertiesT Properties,
+                            Reduction Redu, KernelType KernelFunc) {
+  NDRangeReduction<Strategy>::template run<KernelName>(
+      CGH, Queue, NDRange, Properties, Redu, KernelFunc);
+}
+
 __SYCL_EXPORT uint32_t
 reduGetMaxNumConcurrentWorkGroups(std::shared_ptr<queue_impl> Queue);
 
@@ -2213,161 +2300,66 @@ void reduction_parallel_for(handler &CGH,
       KernelFunc(getDelinearizedId(Range, I), Reducer);
   };
 
-  if constexpr (Reduction::has_fast_reduce)
-    reduCGFuncForRangeFastReduce<KernelName>(CGH, UpdatedKernelFunc, NDRange,
-                                             Properties, Redu);
-  else if constexpr (Reduction::has_fast_atomics)
-    reduCGFuncForRangeFastAtomics<KernelName>(CGH, UpdatedKernelFunc, NDRange,
-                                              Properties, Redu);
-  else
-    reduCGFuncForRangeBasic<KernelName>(CGH, UpdatedKernelFunc, NDRange,
-                                        Properties, Redu);
-}
-
-template <typename KernelName, int Dims, typename PropertiesT,
-          typename KernelType, typename Reduction>
-void reduction_parallel_for_basic_impl(
-    handler &CGH, std::shared_ptr<detail::queue_impl> Queue,
-    nd_range<Dims> Range, PropertiesT Properties, Reduction Redu,
-    KernelType KernelFunc) {
-  // This parallel_for() is lowered to the following sequence:
-  // 1) Call a kernel that a) call user's lambda function and b) performs
-  //    one iteration of reduction, storing the partial reductions/sums
-  //    to either a newly created global buffer or to user's reduction
-  //    accessor. So, if the original 'Range' has totally
-  //    N1 elements and work-group size is W, then after the first iteration
-  //    there will be N2 partial sums where N2 = N1 / W.
-  //    If (N2 == 1) then the partial sum is written to user's accessor.
-  //    Otherwise, a new global buffer is created and partial sums are written
-  //    to it.
-  // 2) Call an aux kernel (if necessary, i.e. if N2 > 1) as many times as
-  //    necessary to reduce all partial sums into one final sum.
-
-  // Before running the kernels, check that device has enough local memory
-  // to hold local arrays that may be required for the reduction algorithm.
-  // TODO: If the work-group-size is limited by the local memory, then
-  // a special version of the main kernel may be created. The one that would
-  // not use local accessors, which means it would not do the reduction in
-  // the main kernel, but simply generate Range.get_global_range.size() number
-  // of partial sums, leaving the reduction work to the additional/aux
-  // kernels.
-  constexpr bool HFR = Reduction::has_fast_reduce;
-  size_t OneElemSize = HFR ? 0 : sizeof(typename Reduction::result_type);
-  // TODO: currently the maximal work group size is determined for the given
-  // queue/device, while it may be safer to use queries to the kernel compiled
-  // for the device.
-  size_t MaxWGSize = reduGetMaxWGSize(Queue, OneElemSize);
-  if (Range.get_local_range().size() > MaxWGSize)
-    throw sycl::runtime_error("The implementation handling parallel_for with"
-                              " reduction requires work group size not bigger"
-                              " than " +
-                                  std::to_string(MaxWGSize),
-                              PI_ERROR_INVALID_WORK_GROUP_SIZE);
-
-  // 1. Call the kernel that includes user's lambda function.
-  // We only call this basic version when we can't use atomics to do the final
-  // reduction.
-  assert(!Reduction::has_fast_atomics);
-  if constexpr (Reduction::has_fast_reduce) {
-    reduCGFuncForNDRangeFastReduceOnly<KernelName, KernelType>(
-        CGH, KernelFunc, Range, Properties, Redu);
-  } else {
-    reduCGFuncForNDRangeBasic<KernelName, KernelType>(CGH, KernelFunc, Range,
-                                                      Properties, Redu);
-  }
-  reduction::finalizeHandler(CGH);
-
-  // 2. Run the additional kernel as many times as needed to reduce
-  // all partial sums into one scalar.
-
-  // TODO: Create a special slow/sequential version of the kernel that would
-  // handle the reduction instead of reporting an assert below.
-  if (MaxWGSize <= 1)
-    throw sycl::runtime_error("The implementation handling parallel_for with "
-                              "reduction requires the maximal work group "
-                              "size to be greater than 1 to converge. "
-                              "The maximal work group size depends on the "
-                              "device and the size of the objects passed to "
-                              "the reduction.",
-                              PI_ERROR_INVALID_WORK_GROUP_SIZE);
-  size_t NWorkItems = Range.get_group_range().size();
-  while (NWorkItems > 1) {
-    reduction::withAuxHandler(CGH, [&](handler &AuxHandler) {
-      NWorkItems = reduAuxCGFunc<KernelName, KernelType>(AuxHandler, NWorkItems,
-                                                         MaxWGSize, Redu);
-    });
-  } // end while (NWorkItems > 1)
-
-  if (Reduction::is_usm) {
-    reduction::withAuxHandler(CGH, [&](handler &CopyHandler) {
-      reduSaveFinalResultToUserMem<KernelName>(CopyHandler, Redu);
-    });
-  }
+  constexpr auto Strategy = [&]() {
+    if constexpr (Reduction::has_fast_reduce)
+      return reduction::strategy::group_reduce_and_last_wg_detection;
+    else if constexpr (Reduction::has_fast_atomics)
+      return reduction::strategy::local_atomic_and_atomic_cross_wg;
+    else
+      return reduction::strategy::range_basic;
+  }();
+
+  reduction_parallel_for<KernelName, Strategy>(CGH, Queue, NDRange, Properties,
+                                               Redu, UpdatedKernelFunc);
 }
 
-template <typename KernelName, int Dims, typename PropertiesT,
-          typename KernelType, typename Reduction>
-void reduction_parallel_for(handler &CGH,
-                            std::shared_ptr<detail::queue_impl> Queue,
-                            nd_range<Dims> Range, PropertiesT Properties,
-                            Reduction Redu, KernelType KernelFunc) {
-  if constexpr (Reduction::has_float64_atomics) {
-    if (detail::getDeviceFromHandler(CGH).has(aspect::atomic64))
-      return reduCGFuncForNDRangeBothFastReduceAndAtomics<KernelName>(
-          CGH, KernelFunc, Range, Properties, Redu);
-
-    return reduction_parallel_for_basic_impl<KernelName>(
-        CGH, Queue, Range, Properties, Redu, KernelFunc);
-  } else if constexpr (Reduction::has_fast_atomics) {
-    if constexpr (Reduction::has_fast_reduce) {
-      return reduCGFuncForNDRangeBothFastReduceAndAtomics<KernelName,
-                                                          KernelType>(
-          CGH, KernelFunc, Range, Properties, Redu);
-    } else {
-      return reduCGFuncForNDRangeFastAtomicsOnly<KernelName, KernelType>(
-          CGH, KernelFunc, Range, Properties, Redu);
-    }
-  } else {
-    return reduction_parallel_for_basic_impl<KernelName>(
-        CGH, Queue, Range, Properties, Redu, KernelFunc);
+template <> struct NDRangeReduction<reduction::strategy::multi> {
+  template <typename KernelName, int Dims, typename PropertiesT,
+            typename... RestT>
+  static void run(handler &CGH, std::shared_ptr<detail::queue_impl> &Queue,
+                  nd_range<Dims> NDRange, PropertiesT &Properties,
+                  RestT... Rest) {
+    std::tuple<RestT...> ArgsTuple(Rest...);
+    constexpr size_t NumArgs = sizeof...(RestT);
+    auto KernelFunc = std::get<NumArgs - 1>(ArgsTuple);
+    auto ReduIndices = std::make_index_sequence<NumArgs - 1>();
+    auto ReduTuple = detail::tuple_select_elements(ArgsTuple, ReduIndices);
+
+    size_t LocalMemPerWorkItem = reduGetMemPerWorkItem(ReduTuple, ReduIndices);
+    // TODO: currently the maximal work group size is determined for the given
+    // queue/device, while it is safer to use queries to the kernel compiled
+    // for the device.
+    size_t MaxWGSize = reduGetMaxWGSize(Queue, LocalMemPerWorkItem);
+    if (NDRange.get_local_range().size() > MaxWGSize)
+      throw sycl::runtime_error("The implementation handling parallel_for with"
+                                " reduction requires work group size not bigger"
+                                " than " +
+                                    std::to_string(MaxWGSize),
+                                PI_ERROR_INVALID_WORK_GROUP_SIZE);
+
+    reduCGFuncMulti<KernelName>(CGH, KernelFunc, NDRange, Properties, ReduTuple,
+                                ReduIndices);
+    reduction::finalizeHandler(CGH);
+
+    size_t NWorkItems = NDRange.get_group_range().size();
+    while (NWorkItems > 1) {
+      reduction::withAuxHandler(CGH, [&](handler &AuxHandler) {
+        NWorkItems = reduAuxCGFunc<KernelName, decltype(KernelFunc)>(
+            AuxHandler, NWorkItems, MaxWGSize, ReduTuple, ReduIndices);
+      });
+    } // end while (NWorkItems > 1)
   }
-}
+};
 
 template <typename KernelName, int Dims, typename PropertiesT,
           typename... RestT>
 void reduction_parallel_for(handler &CGH,
                             std::shared_ptr<detail::queue_impl> Queue,
-                            nd_range<Dims> Range, PropertiesT Properties,
+                            nd_range<Dims> NDRange, PropertiesT Properties,
                             RestT... Rest) {
-  std::tuple<RestT...> ArgsTuple(Rest...);
-  constexpr size_t NumArgs = sizeof...(RestT);
-  auto KernelFunc = std::get<NumArgs - 1>(ArgsTuple);
-  auto ReduIndices = std::make_index_sequence<NumArgs - 1>();
-  auto ReduTuple = detail::tuple_select_elements(ArgsTuple, ReduIndices);
-
-  size_t LocalMemPerWorkItem = reduGetMemPerWorkItem(ReduTuple, ReduIndices);
-  // TODO: currently the maximal work group size is determined for the given
-  // queue/device, while it is safer to use queries to the kernel compiled
-  // for the device.
-  size_t MaxWGSize = reduGetMaxWGSize(Queue, LocalMemPerWorkItem);
-  if (Range.get_local_range().size() > MaxWGSize)
-    throw sycl::runtime_error("The implementation handling parallel_for with"
-                              " reduction requires work group size not bigger"
-                              " than " +
-                                  std::to_string(MaxWGSize),
-                              PI_ERROR_INVALID_WORK_GROUP_SIZE);
-
-  reduCGFuncMulti<KernelName>(CGH, KernelFunc, Range, Properties, ReduTuple,
-                              ReduIndices);
-  reduction::finalizeHandler(CGH);
-
-  size_t NWorkItems = Range.get_group_range().size();
-  while (NWorkItems > 1) {
-    reduction::withAuxHandler(CGH, [&](handler &AuxHandler) {
-      NWorkItems = reduAuxCGFunc<KernelName, decltype(KernelFunc)>(
-          AuxHandler, NWorkItems, MaxWGSize, ReduTuple, ReduIndices);
-    });
-  } // end while (NWorkItems > 1)
+  constexpr auto Strategy = reduction::strategy::multi;
+  NDRangeReduction<Strategy>::template run<KernelName>(CGH, Queue, NDRange,
+                                                       Properties, Rest...);
 }
 } // namespace detail
 
diff --git a/sycl/include/sycl/reduction_fwd.hpp b/sycl/include/sycl/reduction_fwd.hpp
new file mode 100644
index 0000000000000..a7a265f922ed2
--- /dev/null
+++ b/sycl/include/sycl/reduction_fwd.hpp
@@ -0,0 +1,67 @@
+#pragma once
+
+// To be included in <sycl/handler.hpp>. Not that reductions implementation need
+// complete sycl::handler type so we cannot include whole <sycl/reduction.hpp>
+// there.
+
+#include <sycl/detail/common.hpp>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace detail {
+template <typename T, class BinaryOperation, int Dims, size_t Extent,
+          typename RedOutVar>
+class reduction_impl_algo;
+
+namespace reduction {
+enum class strategy : int {
+  auto_select,
+
+  // These three are only auto-selected for sycl::range entry point.
+  group_reduce_and_last_wg_detection,
+  local_atomic_and_atomic_cross_wg,
+  range_basic,
+
+  group_reduce_and_atomic_cross_wg,
+  local_mem_tree_and_atomic_cross_wg,
+  group_reduce_and_multiple_kernels,
+  basic,
+
+  multi,
+};
+
+// Reductions implementation need access to private members of handler. Those
+// are limited to those below.
+inline void finalizeHandler(handler &CGH);
+template <class FunctorTy> void withAuxHandler(handler &CGH, FunctorTy Func);
+} // namespace reduction
+
+template <typename KernelName, int Dims, typename PropertiesT,
+          typename KernelType, typename Reduction>
+void reduction_parallel_for(handler &CGH,
+                            std::shared_ptr<detail::queue_impl> Queue,
+                            range<Dims> Range, PropertiesT Properties,
+                            Reduction Redu, KernelType KernelFunc);
+
+template <typename KernelName,
+          reduction::strategy Strategy = reduction::strategy::auto_select,
+          int Dims, typename PropertiesT, typename KernelType,
+          typename Reduction>
+void reduction_parallel_for(handler &CGH,
+                            std::shared_ptr<detail::queue_impl> Queue,
+                            nd_range<Dims> NDRange, PropertiesT Properties,
+                            Reduction Redu, KernelType KernelFunc);
+
+template <typename KernelName, int Dims, typename PropertiesT,
+          typename... RestT>
+void reduction_parallel_for(handler &CGH,
+                            std::shared_ptr<detail::queue_impl> Queue,
+                            nd_range<Dims> NDRange, PropertiesT Properties,
+                            RestT... Rest);
+
+template <typename T> struct IsReduction;
+template <typename FirstT, typename... RestT> struct AreAllButLastReductions;
+
+} // namespace detail
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl
diff --git a/sycl/test/extensions/properties/properties_kernel_device_has.cpp b/sycl/test/extensions/properties/properties_kernel_device_has.cpp
index e25a910ba5cfb..e0f1e932b38a4 100644
--- a/sycl/test/extensions/properties/properties_kernel_device_has.cpp
+++ b/sycl/test/extensions/properties/properties_kernel_device_has.cpp
@@ -51,12 +51,12 @@ int main() {
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel5(){{.*}} #[[DHAttr2]]
   Q.parallel_for<class WGSizeKernel5>(R1, {Ev}, Props, [](id<1>) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel6{{.*}}{{.*}} #[[DHAttr3:[0-9]+]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel6{{.*}}{{.*}} #[[DHAttr3:[0-9]+]]
   Q.parallel_for<class WGSizeKernel6>(R1, Props, Redu1, [](id<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel7{{.*}}{{.*}} #[[DHAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel7{{.*}}{{.*}} #[[DHAttr3]]
   Q.parallel_for<class WGSizeKernel7>(R1, Ev, Props, Redu1,
                                       [](id<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel8{{.*}}{{.*}} #[[DHAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel8{{.*}}{{.*}} #[[DHAttr3]]
   Q.parallel_for<class WGSizeKernel8>(R1, {Ev}, Props, Redu1,
                                       [](id<1>, auto &) {});
 
@@ -67,23 +67,23 @@ int main() {
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel11(){{.*}} #[[DHAttr2]]
   Q.parallel_for<class WGSizeKernel11>(NDR1, {Ev}, Props, [](nd_item<1>) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel12{{.*}}{{.*}} #[[DHAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel12{{.*}}{{.*}} #[[DHAttr3]]
   Q.parallel_for<class WGSizeKernel12>(NDR1, Props, Redu1,
                                        [](nd_item<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel13{{.*}}{{.*}} #[[DHAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel13{{.*}}{{.*}} #[[DHAttr3]]
   Q.parallel_for<class WGSizeKernel13>(NDR1, Ev, Props, Redu1,
                                        [](nd_item<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel14{{.*}}{{.*}} #[[DHAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel14{{.*}}{{.*}} #[[DHAttr3]]
   Q.parallel_for<class WGSizeKernel14>(NDR1, {Ev}, Props, Redu1,
                                        [](nd_item<1>, auto &) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel15{{.*}}{{.*}} #[[DHAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel15{{.*}}{{.*}} #[[DHAttr3]]
   Q.parallel_for<class WGSizeKernel15>(NDR1, Props, Redu1, Redu2,
                                        [](nd_item<1>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel16{{.*}}{{.*}} #[[DHAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel16{{.*}}{{.*}} #[[DHAttr3]]
   Q.parallel_for<class WGSizeKernel16>(NDR1, Ev, Props, Redu1, Redu2,
                                        [](nd_item<1>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel17{{.*}}{{.*}} #[[DHAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel17{{.*}}{{.*}} #[[DHAttr3]]
   Q.parallel_for<class WGSizeKernel17>(NDR1, {Ev}, Props, Redu1, Redu2,
                                        [](nd_item<1>, auto &, auto &) {});
 
@@ -97,7 +97,7 @@ int main() {
     CGH.parallel_for<class WGSizeKernel19>(R1, Props, [](id<1>) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel20{{.*}}{{.*}} #[[DHAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel20{{.*}}{{.*}} #[[DHAttr3]]
   Q.submit([&](handler &CGH) {
     CGH.parallel_for<class WGSizeKernel20>(R1, Props, Redu1,
                                            [](id<1>, auto &) {});
@@ -108,13 +108,13 @@ int main() {
     CGH.parallel_for<class WGSizeKernel21>(NDR1, Props, [](nd_item<1>) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel22{{.*}}{{.*}} #[[DHAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel22{{.*}}{{.*}} #[[DHAttr3]]
   Q.submit([&](handler &CGH) {
     CGH.parallel_for<class WGSizeKernel22>(NDR1, Props, Redu1,
                                            [](nd_item<1>, auto &) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel23{{.*}}{{.*}} #[[DHAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel23{{.*}}{{.*}} #[[DHAttr3]]
   Q.submit([&](handler &CGH) {
     CGH.parallel_for<class WGSizeKernel23>(NDR1, Props, Redu1, Redu2,
                                            [](nd_item<1>, auto &, auto &) {});
diff --git a/sycl/test/extensions/properties/properties_kernel_sub_group_size.cpp b/sycl/test/extensions/properties/properties_kernel_sub_group_size.cpp
index 96fbb460788db..80cc373aa066f 100644
--- a/sycl/test/extensions/properties/properties_kernel_sub_group_size.cpp
+++ b/sycl/test/extensions/properties/properties_kernel_sub_group_size.cpp
@@ -48,31 +48,31 @@ int main() {
   // CHECK-IR: spir_kernel void @{{.*}}SGSizeKernel17(){{.*}} #[[SGSizeAttr2]]
   Q.parallel_for<class SGSizeKernel17>(R3, {Ev}, Props, [](sycl::id<3>) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel18{{.*}}{{.*}} #[[SGSizeAttr3:[0-9]+]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel18{{.*}}{{.*}} #[[SGSizeAttr3:[0-9]+]]
   Q.parallel_for<class SGSizeKernel18>(R1, Props, Redu1,
                                        [](sycl::id<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel19{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel19{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel19>(R1, Ev, Props, Redu1,
                                        [](sycl::id<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel20{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel20{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel20>(R1, {Ev}, Props, Redu1,
                                        [](sycl::id<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel21{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel21{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel21>(R2, Props, Redu1,
                                        [](sycl::id<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel22{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel22{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel22>(R2, Ev, Props, Redu1,
                                        [](sycl::id<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel23{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel23{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel23>(R2, {Ev}, Props, Redu1,
                                        [](sycl::id<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel24{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel24{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel24>(R3, Props, Redu1,
                                        [](sycl::id<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel25{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel25{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel25>(R3, Ev, Props, Redu1,
                                        [](sycl::id<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel26{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel26{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel26>(R3, {Ev}, Props, Redu1,
                                        [](sycl::id<3>, auto &) {});
 
@@ -101,59 +101,59 @@ int main() {
   Q.parallel_for<class SGSizeKernel35>(NDR3, {Ev}, Props,
                                        [](sycl::nd_item<3>) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel36{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel36{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel36>(NDR1, Props, Redu1,
                                        [](sycl::nd_item<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel37{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel37{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel37>(NDR1, Ev, Props, Redu1,
                                        [](sycl::nd_item<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel38{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel38{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel38>(NDR1, {Ev}, Props, Redu1,
                                        [](sycl::nd_item<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel39{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel39{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel39>(NDR2, Props, Redu1,
                                        [](sycl::nd_item<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel40{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel40{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel40>(NDR2, Ev, Props, Redu1,
                                        [](sycl::nd_item<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel41{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel41{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel41>(NDR2, {Ev}, Props, Redu1,
                                        [](sycl::nd_item<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel42{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel42{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel42>(NDR3, Props, Redu1,
                                        [](sycl::nd_item<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel43{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel43{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel43>(NDR3, Ev, Props, Redu1,
                                        [](sycl::nd_item<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel44{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel44{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel44>(NDR3, {Ev}, Props, Redu1,
                                        [](sycl::nd_item<3>, auto &) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel45{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel45{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel45>(NDR1, Props, Redu1, Redu2,
                                        [](sycl::nd_item<1>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel46{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel46{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel46>(NDR1, Ev, Props, Redu1, Redu2,
                                        [](sycl::nd_item<1>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel47{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel47{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel47>(NDR1, {Ev}, Props, Redu1, Redu2,
                                        [](sycl::nd_item<1>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel48{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel48{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel48>(NDR2, Props, Redu1, Redu2,
                                        [](sycl::nd_item<2>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel49{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel49{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel49>(NDR2, Ev, Props, Redu1, Redu2,
                                        [](sycl::nd_item<2>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel50{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel50{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel50>(NDR2, {Ev}, Props, Redu1, Redu2,
                                        [](sycl::nd_item<2>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel51{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel51{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel51>(NDR3, Props, Redu1, Redu2,
                                        [](sycl::nd_item<3>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel52{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel52{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel52>(NDR3, Ev, Props, Redu1, Redu2,
                                        [](sycl::nd_item<3>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel53{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel53{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.parallel_for<class SGSizeKernel53>(NDR3, {Ev}, Props, Redu1, Redu2,
                                        [](sycl::nd_item<3>, auto &, auto &) {});
 
@@ -183,17 +183,17 @@ int main() {
     CGH.parallel_for<class SGSizeKernel59>(R3, Props, [](sycl::id<3>) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel60{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel60{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class SGSizeKernel60>(R1, Props, Redu1,
                                            [](sycl::id<1>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel61{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel61{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class SGSizeKernel61>(R2, Props, Redu1,
                                            [](sycl::id<2>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel62{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel62{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class SGSizeKernel62>(R3, Props, Redu1,
                                            [](sycl::id<3>, auto &) {});
@@ -215,33 +215,33 @@ int main() {
                                            [](sycl::nd_item<3>) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel66{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel66{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class SGSizeKernel66>(NDR1, Props, Redu1,
                                            [](sycl::nd_item<1>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel67{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel67{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class SGSizeKernel67>(NDR2, Props, Redu1,
                                            [](sycl::nd_item<2>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel68{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel68{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class SGSizeKernel68>(NDR3, Props, Redu1,
                                            [](sycl::nd_item<3>, auto &) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel69{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel69{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class SGSizeKernel69>(
         NDR1, Props, Redu1, Redu2, [](sycl::nd_item<1>, auto &, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel70{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel70{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class SGSizeKernel70>(
         NDR2, Props, Redu1, Redu2, [](sycl::nd_item<2>, auto &, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}SGSizeKernel71{{.*}}{{.*}} #[[SGSizeAttr3]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}SGSizeKernel71{{.*}}{{.*}} #[[SGSizeAttr3]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class SGSizeKernel71>(
         NDR3, Props, Redu1, Redu2, [](sycl::nd_item<3>, auto &, auto &) {});
diff --git a/sycl/test/extensions/properties/properties_kernel_work_group_size.cpp b/sycl/test/extensions/properties/properties_kernel_work_group_size.cpp
index 336b0825cfc19..7f66bed0fb9b9 100644
--- a/sycl/test/extensions/properties/properties_kernel_work_group_size.cpp
+++ b/sycl/test/extensions/properties/properties_kernel_work_group_size.cpp
@@ -64,31 +64,31 @@ int main() {
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel17(){{.*}} #[[WGSizeAttr6]]
   Q.parallel_for<class WGSizeKernel17>(R3, {Ev}, Props3, [](sycl::id<3>) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel18{{.*}}{{.*}} #[[WGSizeAttr7:[0-9]+]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel18{{.*}}{{.*}} #[[WGSizeAttr7:[0-9]+]]
   Q.parallel_for<class WGSizeKernel18>(R1, Props1, Redu1,
                                        [](sycl::id<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel19{{.*}}{{.*}} #[[WGSizeAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel19{{.*}}{{.*}} #[[WGSizeAttr7]]
   Q.parallel_for<class WGSizeKernel19>(R1, Ev, Props1, Redu1,
                                        [](sycl::id<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel20{{.*}}{{.*}} #[[WGSizeAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel20{{.*}}{{.*}} #[[WGSizeAttr7]]
   Q.parallel_for<class WGSizeKernel20>(R1, {Ev}, Props1, Redu1,
                                        [](sycl::id<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel21{{.*}}{{.*}} #[[WGSizeAttr8:[0-9]+]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel21{{.*}}{{.*}} #[[WGSizeAttr8:[0-9]+]]
   Q.parallel_for<class WGSizeKernel21>(R2, Props2, Redu1,
                                        [](sycl::id<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel22{{.*}}{{.*}} #[[WGSizeAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel22{{.*}}{{.*}} #[[WGSizeAttr8]]
   Q.parallel_for<class WGSizeKernel22>(R2, Ev, Props2, Redu1,
                                        [](sycl::id<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel23{{.*}}{{.*}} #[[WGSizeAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel23{{.*}}{{.*}} #[[WGSizeAttr8]]
   Q.parallel_for<class WGSizeKernel23>(R2, {Ev}, Props2, Redu1,
                                        [](sycl::id<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel24{{.*}}{{.*}} #[[WGSizeAttr9:[0-9]+]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel24{{.*}}{{.*}} #[[WGSizeAttr9:[0-9]+]]
   Q.parallel_for<class WGSizeKernel24>(R3, Props3, Redu1,
                                        [](sycl::id<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel25{{.*}}{{.*}} #[[WGSizeAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel25{{.*}}{{.*}} #[[WGSizeAttr9]]
   Q.parallel_for<class WGSizeKernel25>(R3, Ev, Props3, Redu1,
                                        [](sycl::id<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel26{{.*}}{{.*}} #[[WGSizeAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel26{{.*}}{{.*}} #[[WGSizeAttr9]]
   Q.parallel_for<class WGSizeKernel26>(R3, {Ev}, Props3, Redu1,
                                        [](sycl::id<3>, auto &) {});
 
@@ -117,59 +117,59 @@ int main() {
   Q.parallel_for<class WGSizeKernel35>(NDR3, {Ev}, Props3,
                                        [](sycl::nd_item<3>) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel36{{.*}}{{.*}} #[[WGSizeAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel36{{.*}}{{.*}} #[[WGSizeAttr7]]
   Q.parallel_for<class WGSizeKernel36>(NDR1, Props1, Redu1,
                                        [](sycl::nd_item<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel37{{.*}}{{.*}} #[[WGSizeAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel37{{.*}}{{.*}} #[[WGSizeAttr7]]
   Q.parallel_for<class WGSizeKernel37>(NDR1, Ev, Props1, Redu1,
                                        [](sycl::nd_item<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel38{{.*}}{{.*}} #[[WGSizeAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel38{{.*}}{{.*}} #[[WGSizeAttr7]]
   Q.parallel_for<class WGSizeKernel38>(NDR1, {Ev}, Props1, Redu1,
                                        [](sycl::nd_item<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel39{{.*}}{{.*}} #[[WGSizeAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel39{{.*}}{{.*}} #[[WGSizeAttr8]]
   Q.parallel_for<class WGSizeKernel39>(NDR2, Props2, Redu1,
                                        [](sycl::nd_item<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel40{{.*}}{{.*}} #[[WGSizeAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel40{{.*}}{{.*}} #[[WGSizeAttr8]]
   Q.parallel_for<class WGSizeKernel40>(NDR2, Ev, Props2, Redu1,
                                        [](sycl::nd_item<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel41{{.*}}{{.*}} #[[WGSizeAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel41{{.*}}{{.*}} #[[WGSizeAttr8]]
   Q.parallel_for<class WGSizeKernel41>(NDR2, {Ev}, Props2, Redu1,
                                        [](sycl::nd_item<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel42{{.*}}{{.*}} #[[WGSizeAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel42{{.*}}{{.*}} #[[WGSizeAttr9]]
   Q.parallel_for<class WGSizeKernel42>(NDR3, Props3, Redu1,
                                        [](sycl::nd_item<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel43{{.*}}{{.*}} #[[WGSizeAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel43{{.*}}{{.*}} #[[WGSizeAttr9]]
   Q.parallel_for<class WGSizeKernel43>(NDR3, Ev, Props3, Redu1,
                                        [](sycl::nd_item<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel44{{.*}}{{.*}} #[[WGSizeAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel44{{.*}}{{.*}} #[[WGSizeAttr9]]
   Q.parallel_for<class WGSizeKernel44>(NDR3, {Ev}, Props3, Redu1,
                                        [](sycl::nd_item<3>, auto &) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel45{{.*}}{{.*}} #[[WGSizeAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel45{{.*}}{{.*}} #[[WGSizeAttr7]]
   Q.parallel_for<class WGSizeKernel45>(NDR1, Props1, Redu1, Redu2,
                                        [](sycl::nd_item<1>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel46{{.*}}{{.*}} #[[WGSizeAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel46{{.*}}{{.*}} #[[WGSizeAttr7]]
   Q.parallel_for<class WGSizeKernel46>(NDR1, Ev, Props1, Redu1, Redu2,
                                        [](sycl::nd_item<1>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel47{{.*}}{{.*}} #[[WGSizeAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel47{{.*}}{{.*}} #[[WGSizeAttr7]]
   Q.parallel_for<class WGSizeKernel47>(NDR1, {Ev}, Props1, Redu1, Redu2,
                                        [](sycl::nd_item<1>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel48{{.*}}{{.*}} #[[WGSizeAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel48{{.*}}{{.*}} #[[WGSizeAttr8]]
   Q.parallel_for<class WGSizeKernel48>(NDR2, Props2, Redu1, Redu2,
                                        [](sycl::nd_item<2>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel49{{.*}}{{.*}} #[[WGSizeAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel49{{.*}}{{.*}} #[[WGSizeAttr8]]
   Q.parallel_for<class WGSizeKernel49>(NDR2, Ev, Props2, Redu1, Redu2,
                                        [](sycl::nd_item<2>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel50{{.*}}{{.*}} #[[WGSizeAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel50{{.*}}{{.*}} #[[WGSizeAttr8]]
   Q.parallel_for<class WGSizeKernel50>(NDR2, {Ev}, Props2, Redu1, Redu2,
                                        [](sycl::nd_item<2>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel51{{.*}}{{.*}} #[[WGSizeAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel51{{.*}}{{.*}} #[[WGSizeAttr9]]
   Q.parallel_for<class WGSizeKernel51>(NDR3, Props3, Redu1, Redu2,
                                        [](sycl::nd_item<3>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel52{{.*}}{{.*}} #[[WGSizeAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel52{{.*}}{{.*}} #[[WGSizeAttr9]]
   Q.parallel_for<class WGSizeKernel52>(NDR3, Ev, Props3, Redu1, Redu2,
                                        [](sycl::nd_item<3>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel53{{.*}}{{.*}} #[[WGSizeAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel53{{.*}}{{.*}} #[[WGSizeAttr9]]
   Q.parallel_for<class WGSizeKernel53>(NDR3, {Ev}, Props3, Redu1, Redu2,
                                        [](sycl::nd_item<3>, auto &, auto &) {});
 
@@ -199,17 +199,17 @@ int main() {
     CGH.parallel_for<class WGSizeKernel59>(R3, Props3, [](sycl::id<3>) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel60{{.*}}{{.*}} #[[WGSizeAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel60{{.*}}{{.*}} #[[WGSizeAttr7]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeKernel60>(R1, Props1, Redu1,
                                            [](sycl::id<1>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel61{{.*}}{{.*}} #[[WGSizeAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel61{{.*}}{{.*}} #[[WGSizeAttr8]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeKernel61>(R2, Props2, Redu1,
                                            [](sycl::id<2>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel62{{.*}}{{.*}} #[[WGSizeAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel62{{.*}}{{.*}} #[[WGSizeAttr9]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeKernel62>(R3, Props3, Redu1,
                                            [](sycl::id<3>, auto &) {});
@@ -231,33 +231,33 @@ int main() {
                                            [](sycl::nd_item<3>) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel66{{.*}}{{.*}} #[[WGSizeAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel66{{.*}}{{.*}} #[[WGSizeAttr7]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeKernel66>(NDR1, Props1, Redu1,
                                            [](sycl::nd_item<1>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel67{{.*}}{{.*}} #[[WGSizeAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel67{{.*}}{{.*}} #[[WGSizeAttr8]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeKernel67>(NDR2, Props2, Redu1,
                                            [](sycl::nd_item<2>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel68{{.*}}{{.*}} #[[WGSizeAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel68{{.*}}{{.*}} #[[WGSizeAttr9]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeKernel68>(NDR3, Props3, Redu1,
                                            [](sycl::nd_item<3>, auto &) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel69{{.*}}{{.*}} #[[WGSizeAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel69{{.*}}{{.*}} #[[WGSizeAttr7]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeKernel69>(
         NDR1, Props1, Redu1, Redu2, [](sycl::nd_item<1>, auto &, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel70{{.*}}{{.*}} #[[WGSizeAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel70{{.*}}{{.*}} #[[WGSizeAttr8]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeKernel70>(
         NDR2, Props2, Redu1, Redu2, [](sycl::nd_item<2>, auto &, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeKernel71{{.*}}{{.*}} #[[WGSizeAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel71{{.*}}{{.*}} #[[WGSizeAttr9]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeKernel71>(
         NDR3, Props3, Redu1, Redu2, [](sycl::nd_item<3>, auto &, auto &) {});
diff --git a/sycl/test/extensions/properties/properties_kernel_work_group_size_hint.cpp b/sycl/test/extensions/properties/properties_kernel_work_group_size_hint.cpp
index 2dd3d5e8494fe..a844b484b8b51 100644
--- a/sycl/test/extensions/properties/properties_kernel_work_group_size_hint.cpp
+++ b/sycl/test/extensions/properties/properties_kernel_work_group_size_hint.cpp
@@ -67,31 +67,31 @@ int main() {
   Q.parallel_for<class WGSizeHintKernel17>(R3, {Ev}, Props3,
                                            [](sycl::id<3>) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel18{{.*}}{{.*}} #[[WGSizeHintAttr7:[0-9]+]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel18{{.*}}{{.*}} #[[WGSizeHintAttr7:[0-9]+]]
   Q.parallel_for<class WGSizeHintKernel18>(R1, Props1, Redu1,
                                            [](sycl::id<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel19{{.*}}{{.*}} #[[WGSizeHintAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel19{{.*}}{{.*}} #[[WGSizeHintAttr7]]
   Q.parallel_for<class WGSizeHintKernel19>(R1, Ev, Props1, Redu1,
                                            [](sycl::id<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel20{{.*}}{{.*}} #[[WGSizeHintAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel20{{.*}}{{.*}} #[[WGSizeHintAttr7]]
   Q.parallel_for<class WGSizeHintKernel20>(R1, {Ev}, Props1, Redu1,
                                            [](sycl::id<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel21{{.*}}{{.*}} #[[WGSizeHintAttr8:[0-9]+]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel21{{.*}}{{.*}} #[[WGSizeHintAttr8:[0-9]+]]
   Q.parallel_for<class WGSizeHintKernel21>(R2, Props2, Redu1,
                                            [](sycl::id<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel22{{.*}}{{.*}} #[[WGSizeHintAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel22{{.*}}{{.*}} #[[WGSizeHintAttr8]]
   Q.parallel_for<class WGSizeHintKernel22>(R2, Ev, Props2, Redu1,
                                            [](sycl::id<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel23{{.*}}{{.*}} #[[WGSizeHintAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel23{{.*}}{{.*}} #[[WGSizeHintAttr8]]
   Q.parallel_for<class WGSizeHintKernel23>(R2, {Ev}, Props2, Redu1,
                                            [](sycl::id<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel24{{.*}}{{.*}} #[[WGSizeHintAttr9:[0-9]+]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel24{{.*}}{{.*}} #[[WGSizeHintAttr9:[0-9]+]]
   Q.parallel_for<class WGSizeHintKernel24>(R3, Props3, Redu1,
                                            [](sycl::id<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel25{{.*}}{{.*}} #[[WGSizeHintAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel25{{.*}}{{.*}} #[[WGSizeHintAttr9]]
   Q.parallel_for<class WGSizeHintKernel25>(R3, Ev, Props3, Redu1,
                                            [](sycl::id<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel26{{.*}}{{.*}} #[[WGSizeHintAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel26{{.*}}{{.*}} #[[WGSizeHintAttr9]]
   Q.parallel_for<class WGSizeHintKernel26>(R3, {Ev}, Props3, Redu1,
                                            [](sycl::id<3>, auto &) {});
 
@@ -123,61 +123,61 @@ int main() {
   Q.parallel_for<class WGSizeHintKernel35>(NDR3, {Ev}, Props3,
                                            [](sycl::nd_item<3>) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel36{{.*}}{{.*}} #[[WGSizeHintAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel36{{.*}}{{.*}} #[[WGSizeHintAttr7]]
   Q.parallel_for<class WGSizeHintKernel36>(NDR1, Props1, Redu1,
                                            [](sycl::nd_item<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel37{{.*}}{{.*}} #[[WGSizeHintAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel37{{.*}}{{.*}} #[[WGSizeHintAttr7]]
   Q.parallel_for<class WGSizeHintKernel37>(NDR1, Ev, Props1, Redu1,
                                            [](sycl::nd_item<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel38{{.*}}{{.*}} #[[WGSizeHintAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel38{{.*}}{{.*}} #[[WGSizeHintAttr7]]
   Q.parallel_for<class WGSizeHintKernel38>(NDR1, {Ev}, Props1, Redu1,
                                            [](sycl::nd_item<1>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel39{{.*}}{{.*}} #[[WGSizeHintAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel39{{.*}}{{.*}} #[[WGSizeHintAttr8]]
   Q.parallel_for<class WGSizeHintKernel39>(NDR2, Props2, Redu1,
                                            [](sycl::nd_item<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel40{{.*}}{{.*}} #[[WGSizeHintAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel40{{.*}}{{.*}} #[[WGSizeHintAttr8]]
   Q.parallel_for<class WGSizeHintKernel40>(NDR2, Ev, Props2, Redu1,
                                            [](sycl::nd_item<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel41{{.*}}{{.*}} #[[WGSizeHintAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel41{{.*}}{{.*}} #[[WGSizeHintAttr8]]
   Q.parallel_for<class WGSizeHintKernel41>(NDR2, {Ev}, Props2, Redu1,
                                            [](sycl::nd_item<2>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel42{{.*}}{{.*}} #[[WGSizeHintAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel42{{.*}}{{.*}} #[[WGSizeHintAttr9]]
   Q.parallel_for<class WGSizeHintKernel42>(NDR3, Props3, Redu1,
                                            [](sycl::nd_item<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel43{{.*}}{{.*}} #[[WGSizeHintAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel43{{.*}}{{.*}} #[[WGSizeHintAttr9]]
   Q.parallel_for<class WGSizeHintKernel43>(NDR3, Ev, Props3, Redu1,
                                            [](sycl::nd_item<3>, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel44{{.*}}{{.*}} #[[WGSizeHintAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel44{{.*}}{{.*}} #[[WGSizeHintAttr9]]
   Q.parallel_for<class WGSizeHintKernel44>(NDR3, {Ev}, Props3, Redu1,
                                            [](sycl::nd_item<3>, auto &) {});
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel45{{.*}}{{.*}} #[[WGSizeHintAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel45{{.*}}{{.*}} #[[WGSizeHintAttr7]]
   Q.parallel_for<class WGSizeHintKernel45>(
       NDR1, Props1, Redu1, Redu2, [](sycl::nd_item<1>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel46{{.*}}{{.*}} #[[WGSizeHintAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel46{{.*}}{{.*}} #[[WGSizeHintAttr7]]
   Q.parallel_for<class WGSizeHintKernel46>(
       NDR1, Ev, Props1, Redu1, Redu2, [](sycl::nd_item<1>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel47{{.*}}{{.*}} #[[WGSizeHintAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel47{{.*}}{{.*}} #[[WGSizeHintAttr7]]
   Q.parallel_for<class WGSizeHintKernel47>(
       NDR1, {Ev}, Props1, Redu1, Redu2,
       [](sycl::nd_item<1>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel48{{.*}}{{.*}} #[[WGSizeHintAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel48{{.*}}{{.*}} #[[WGSizeHintAttr8]]
   Q.parallel_for<class WGSizeHintKernel48>(
       NDR2, Props2, Redu1, Redu2, [](sycl::nd_item<2>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel49{{.*}}{{.*}} #[[WGSizeHintAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel49{{.*}}{{.*}} #[[WGSizeHintAttr8]]
   Q.parallel_for<class WGSizeHintKernel49>(
       NDR2, Ev, Props2, Redu1, Redu2, [](sycl::nd_item<2>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel50{{.*}}{{.*}} #[[WGSizeHintAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel50{{.*}}{{.*}} #[[WGSizeHintAttr8]]
   Q.parallel_for<class WGSizeHintKernel50>(
       NDR2, {Ev}, Props2, Redu1, Redu2,
       [](sycl::nd_item<2>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel51{{.*}}{{.*}} #[[WGSizeHintAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel51{{.*}}{{.*}} #[[WGSizeHintAttr9]]
   Q.parallel_for<class WGSizeHintKernel51>(
       NDR3, Props3, Redu1, Redu2, [](sycl::nd_item<3>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel52{{.*}}{{.*}} #[[WGSizeHintAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel52{{.*}}{{.*}} #[[WGSizeHintAttr9]]
   Q.parallel_for<class WGSizeHintKernel52>(
       NDR3, Ev, Props3, Redu1, Redu2, [](sycl::nd_item<3>, auto &, auto &) {});
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel53{{.*}}{{.*}} #[[WGSizeHintAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel53{{.*}}{{.*}} #[[WGSizeHintAttr9]]
   Q.parallel_for<class WGSizeHintKernel53>(
       NDR3, {Ev}, Props3, Redu1, Redu2,
       [](sycl::nd_item<3>, auto &, auto &) {});
@@ -208,17 +208,17 @@ int main() {
     CGH.parallel_for<class WGSizeHintKernel59>(R3, Props3, [](sycl::id<3>) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel60{{.*}}{{.*}} #[[WGSizeHintAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel60{{.*}}{{.*}} #[[WGSizeHintAttr7]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeHintKernel60>(R1, Props1, Redu1,
                                                [](sycl::id<1>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel61{{.*}}{{.*}} #[[WGSizeHintAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel61{{.*}}{{.*}} #[[WGSizeHintAttr8]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeHintKernel61>(R2, Props2, Redu1,
                                                [](sycl::id<2>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel62{{.*}}{{.*}} #[[WGSizeHintAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel62{{.*}}{{.*}} #[[WGSizeHintAttr9]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeHintKernel62>(R3, Props3, Redu1,
                                                [](sycl::id<3>, auto &) {});
@@ -240,33 +240,33 @@ int main() {
                                                [](sycl::nd_item<3>) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel66{{.*}}{{.*}} #[[WGSizeHintAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel66{{.*}}{{.*}} #[[WGSizeHintAttr7]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeHintKernel66>(NDR1, Props1, Redu1,
                                                [](sycl::nd_item<1>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel67{{.*}}{{.*}} #[[WGSizeHintAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel67{{.*}}{{.*}} #[[WGSizeHintAttr8]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeHintKernel67>(NDR2, Props2, Redu1,
                                                [](sycl::nd_item<2>, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel68{{.*}}{{.*}} #[[WGSizeHintAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel68{{.*}}{{.*}} #[[WGSizeHintAttr9]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeHintKernel68>(NDR3, Props3, Redu1,
                                                [](sycl::nd_item<3>, auto &) {});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel69{{.*}}{{.*}} #[[WGSizeHintAttr7]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel69{{.*}}{{.*}} #[[WGSizeHintAttr7]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeHintKernel69>(
         NDR1, Props1, Redu1, Redu2, [](sycl::nd_item<1>, auto &, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel70{{.*}}{{.*}} #[[WGSizeHintAttr8]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel70{{.*}}{{.*}} #[[WGSizeHintAttr8]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeHintKernel70>(
         NDR2, Props2, Redu1, Redu2, [](sycl::nd_item<2>, auto &, auto &) {});
   });
-  // CHECK-IR: spir_kernel void @{{.*}}main_krn{{.*}}WGSizeHintKernel71{{.*}}{{.*}} #[[WGSizeHintAttr9]]
+  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeHintKernel71{{.*}}{{.*}} #[[WGSizeHintAttr9]]
   Q.submit([&](sycl::handler &CGH) {
     CGH.parallel_for<class WGSizeHintKernel71>(
         NDR3, Props3, Redu1, Redu2, [](sycl::nd_item<3>, auto &, auto &) {});

From 1734edf3d70355f65ed70e3b78553abfd6ab2fd5 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Wed, 9 Nov 2022 08:09:26 -0800
Subject: [PATCH 2/3] License + typo fix

---
 sycl/include/sycl/reduction_fwd.hpp | 14 +++++++++++---
 1 file changed, 11 insertions(+), 3 deletions(-)

diff --git a/sycl/include/sycl/reduction_fwd.hpp b/sycl/include/sycl/reduction_fwd.hpp
index a7a265f922ed2..798e30b5e1ec9 100644
--- a/sycl/include/sycl/reduction_fwd.hpp
+++ b/sycl/include/sycl/reduction_fwd.hpp
@@ -1,8 +1,16 @@
+//==--------- reduction_fwd.hpp - SYCL reduction fwd decl -------*- C++ -*---==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+// ===--------------------------------------------------------------------=== //
+
 #pragma once
 
-// To be included in <sycl/handler.hpp>. Not that reductions implementation need
-// complete sycl::handler type so we cannot include whole <sycl/reduction.hpp>
-// there.
+// To be included in <sycl/handler.hpp>. Note that reductions implementation
+// need complete sycl::handler type so we cannot include whole
+// <sycl/reduction.hpp> there.
 
 #include <sycl/detail/common.hpp>
 

From f06f00e7de608ee81e9fa8b948ccf0bb3c868f54 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Wed, 9 Nov 2022 12:30:44 -0800
Subject: [PATCH 3/3] Rename header

---
 sycl/include/sycl/handler.hpp                                  | 2 +-
 sycl/include/sycl/reduction.hpp                                | 2 +-
 sycl/include/sycl/{reduction_fwd.hpp => reduction_forward.hpp} | 2 +-
 3 files changed, 3 insertions(+), 3 deletions(-)
 rename sycl/include/sycl/{reduction_fwd.hpp => reduction_forward.hpp} (97%)

diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index c03db5fd23e9a..4fffc31a60306 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -30,7 +30,7 @@
 #include <sycl/nd_item.hpp>
 #include <sycl/nd_range.hpp>
 #include <sycl/property_list.hpp>
-#include <sycl/reduction_fwd.hpp>
+#include <sycl/reduction_forward.hpp>
 #include <sycl/sampler.hpp>
 #include <sycl/stl.hpp>
 
diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index 14bd060d6e54b..bc84422033fbe 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -18,7 +18,7 @@
 #include <sycl/kernel.hpp>
 #include <sycl/known_identity.hpp>
 #include <sycl/properties/reduction_properties.hpp>
-#include <sycl/reduction_fwd.hpp>
+#include <sycl/reduction_forward.hpp>
 #include <sycl/usm.hpp>
 
 #include <tuple>
diff --git a/sycl/include/sycl/reduction_fwd.hpp b/sycl/include/sycl/reduction_forward.hpp
similarity index 97%
rename from sycl/include/sycl/reduction_fwd.hpp
rename to sycl/include/sycl/reduction_forward.hpp
index 798e30b5e1ec9..af6145b8b5362 100644
--- a/sycl/include/sycl/reduction_fwd.hpp
+++ b/sycl/include/sycl/reduction_forward.hpp
@@ -1,4 +1,4 @@
-//==--------- reduction_fwd.hpp - SYCL reduction fwd decl -------*- C++ -*---==//
+//==---- reduction_forward.hpp - SYCL reduction forward decl ---*- C++ -*---==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.