[SYCL][Reduction] Fix issue with multiple buffer reductions (#6680)

reduAuxCGFunc has different output target between the case when number
of work groups is one or more, while the kernel executed is the same. As
such we were re-using the same kernel code trying to adjust its name so
that a single one isn't used two times (#WGs check is a run-time one).

We used accessors' types for that but that approach only worked for USM
case (ptr vs buffer accessor for the output type depending on #WGs). In
case when the original reduction variable used a buffer, both type were
the same resulting in a "definition with same mangled name" error. Use
something that has distinct types to name the kernel.

Author: aelovikov-intel

sycl/include/sycl/reduction.hpp

@@ -2272,7 +2272,7 @@ void reduCGFunc(handler &CGH, KernelType KernelFunc,
 
 namespace reduction {
 namespace aux_krn {
-template <class KernelName, class Accessor> struct Multi;
+template <class KernelName, class Predicate> struct Multi;
 } // namespace aux_krn
 } // namespace reduction
 template <typename KernelName, typename KernelType, typename... Reductions,
@@ -2312,7 +2312,7 @@ size_t reduAuxCGFunc(handler &CGH, size_t NWorkItems, size_t MaxWGSize,
     auto AccReduIndices = filterSequence<Reductions...>(Predicate, ReduIndices);
     associateReduAccsWithHandler(CGH, ReduTuple, AccReduIndices);
     using Name = __sycl_reduction_kernel<reduction::aux_krn::Multi, KernelName,
-                                         decltype(OutAccsTuple)>;
+                                         decltype(Predicate)>;
     // TODO: Opportunity to parallelize across number of elements
     range<1> GlobalRange = {HasUniformWG ? NWorkItems : NWorkGroups * WGSize};
     nd_range<1> Range{GlobalRange, range<1>(WGSize)};