[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.

Author: aelovikov-intel

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.