reduction.hpp File Reference

#include <CL/sycl/accessor.hpp>
#include <CL/sycl/atomic.hpp>
#include <CL/sycl/detail/tuple.hpp>
#include <CL/sycl/handler.hpp>
#include <CL/sycl/kernel.hpp>
#include <CL/sycl/known_identity.hpp>
#include <sycl/ext/oneapi/accessor_property_list.hpp>
#include <sycl/ext/oneapi/group_algorithm.hpp>
#include <tuple>

Include dependency graph for reduction.hpp:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Classes
class	cl::sycl::ext::oneapi::detail::reducer< T, BinaryOperation, Dims, Extent, Algorithm, View, Subst >
	Class that is used to represent objects that are passed to user's lambda functions and representing users' reduction variable. More...
struct	cl::sycl::ext::oneapi::detail::ReducerTraits< Reducer >
	Helper class for accessing reducer-defined types in CRTP May prove to be useful for other things later. More...
struct	cl::sycl::ext::oneapi::detail::ReducerTraits< reducer< T, BinaryOperation, Dims, Extent, Algorithm, View, Subst > >
class	cl::sycl::ext::oneapi::detail::combiner< Reducer >
	Use CRTP to avoid redefining shorthand operators in terms of combine. More...
class	cl::sycl::ext::oneapi::detail::reducer< T, BinaryOperation, Dims, Extent, Algorithm, View, enable_if_t< Dims==0 &&Extent==1 &&View==false &&!IsKnownIdentityOp< T, BinaryOperation >::value > >
	Specialization of the generic class 'reducer'. More...
class	cl::sycl::ext::oneapi::detail::reducer< T, BinaryOperation, Dims, Extent, Algorithm, View, enable_if_t< Dims==0 &&Extent==1 &&View==false &&IsKnownIdentityOp< T, BinaryOperation >::value > >
	Specialization of the generic class 'reducer'. More...
class	cl::sycl::ext::oneapi::detail::reducer< T, BinaryOperation, Dims, Extent, Algorithm, View, enable_if_t< Dims==0 &&View==true > >
	Component of 'reducer' class for array reductions, representing a single element of the span (as returned by the subscript operator). More...
class	cl::sycl::ext::oneapi::detail::reducer< T, BinaryOperation, Dims, Extent, Algorithm, View, enable_if_t< Dims==1 &&View==false &&!IsKnownIdentityOp< T, BinaryOperation >::value > >
	Specialization of 'reducer' class for array reductions exposing the subscript operator. More...
class	cl::sycl::ext::oneapi::detail::reducer< T, BinaryOperation, Dims, Extent, Algorithm, View, enable_if_t< Dims==1 &&View==false &&IsKnownIdentityOp< T, BinaryOperation >::value > >
	Specialization of 'reducer' class for array reductions accepting a span in cases where the identity value is known. More...
class	cl::sycl::ext::oneapi::detail::reduction_impl_base
	Base non-template class which is a base class for all reduction implementation classes. More...
class	cl::sycl::ext::oneapi::detail::reduction_impl_common< T, BinaryOperation >
	Templated class for common functionality of all reduction implementation classes. More...
class	cl::sycl::ext::oneapi::detail::default_reduction_algorithm< IsUSM, IsPlaceholder, AccessorDims >
	Types representing specific reduction algorithms Enables reduction_impl_algo to take additional algorithm-specific templates. More...
class	cl::sycl::ext::oneapi::detail::reduction_impl_algo< T, BinaryOperation, Dims, Extent, Algorithm >
	Templated class for implementations of specific reduction algorithms. More...
class	cl::sycl::ext::oneapi::detail::reduction_impl_algo< T, BinaryOperation, Dims, Extent, default_reduction_algorithm< IsUSM, IsPlaceholder, AccessorDims > >
	Original reduction algorithm is the default. More...
struct	cl::sycl::ext::oneapi::detail::AreAllButLastReductions< FirstT, RestT >
	Predicate returning true if all template type parameters except the last one are reductions. More...
struct	cl::sycl::ext::oneapi::detail::AreAllButLastReductions< T >
	Helper specialization of AreAllButLastReductions for one element only. More...
class	cl::sycl::ext::oneapi::detail::reduction_impl< T, BinaryOperation, Dims, Extent, Algorithm >
	This class encapsulates the reduction variable/accessor, the reduction operator and an optional operator identity. More...
class	cl::sycl::ext::oneapi::detail::__sycl_reduction_main_kernel< T1, B1, B2, T2 >
	These are the forward declaration for the classes that help to create names for additional kernels. More...
class	cl::sycl::ext::oneapi::detail::__sycl_reduction_aux_kernel< T1, B1, B2, T2 >
struct	cl::sycl::ext::oneapi::detail::get_reduction_main_kernel_name_t< Name, Type, B1, B2, T3 >
	Helper structs to get additional kernel name types based on given Name and additional template parameters helping to distinguish kernels. More...
struct	cl::sycl::ext::oneapi::detail::get_reduction_main_kernel_name_t< sycl::detail::auto_name, Type, B1, B2, T3 >
struct	cl::sycl::ext::oneapi::detail::get_reduction_aux_kernel_name_t< Name, Type, B1, B2, T3 >
struct	cl::sycl::ext::oneapi::detail::get_reduction_aux_kernel_name_t< sycl::detail::auto_name, Type, B1, B2, T3 >
struct	cl::sycl::ext::oneapi::detail::IsNonUsmReductionPredicate
struct	cl::sycl::ext::oneapi::detail::IsNonUsmReductionPredicate::Func< T >
struct	cl::sycl::ext::oneapi::detail::EmptyReductionPredicate
struct	cl::sycl::ext::oneapi::detail::EmptyReductionPredicate::Func< T >
struct	cl::sycl::ext::oneapi::detail::FilterElement< Cond, I >
struct	cl::sycl::ext::oneapi::detail::IsScalarReduction
struct	cl::sycl::ext::oneapi::detail::IsScalarReduction::Func< Reduction >
struct	cl::sycl::ext::oneapi::detail::IsArrayReduction
struct	cl::sycl::ext::oneapi::detail::IsArrayReduction::Func< Reduction >
struct	cl::sycl::ext::oneapi::has_known_identity< BinaryOperation, AccumulatorT >
struct	cl::sycl::ext::oneapi::known_identity< BinaryOperation, AccumulatorT >

Namespaces
	cl
	We provide new interfaces for matrix muliply in this patch:
	cl::sycl
	cl::sycl::ext
	cl::sycl::ext::oneapi
	cl::sycl::ext::oneapi::detail

Typedefs
template<typename T , class BinaryOperation >
using	cl::sycl::ext::oneapi::detail::IsReduOptForFastAtomicFetch = bool_constant< sycl::detail::is_sgeninteger< T >::value &&sycl::detail::IsValidAtomicType< T >::value &&(sycl::detail::IsPlus< T, BinaryOperation >::value||sycl::detail::IsMinimum< T, BinaryOperation >::value||sycl::detail::IsMaximum< T, BinaryOperation >::value||sycl::detail::IsBitOR< T, BinaryOperation >::value||sycl::detail::IsBitXOR< T, BinaryOperation >::value||sycl::detail::IsBitAND< T, BinaryOperation >::value)>
template<typename T , class BinaryOperation >
using	cl::sycl::ext::oneapi::detail::IsReduOptForAtomic64Add = bool_constant< sycl::detail::IsPlus< T, BinaryOperation >::value &&sycl::detail::is_sgenfloat< T >::value &&(sizeof(T)==4||sizeof(T)==8)>
template<typename T , class BinaryOperation >
using	cl::sycl::ext::oneapi::detail::IsReduOptForFastReduce = bool_constant<((sycl::detail::is_sgeninteger< T >::value &&(sizeof(T)==4||sizeof(T)==8))||sycl::detail::is_sgenfloat< T >::value) &&(sycl::detail::IsPlus< T, BinaryOperation >::value||sycl::detail::IsMinimum< T, BinaryOperation >::value||sycl::detail::IsMaximum< T, BinaryOperation >::value)>
template<typename... Ts>
using	cl::sycl::ext::oneapi::detail::ReduTupleT = sycl::detail::tuple< Ts... >

Functions
template<typename... Ts>
ReduTupleT< Ts... >	cl::sycl::ext::oneapi::detail::makeReduTupleT (Ts... Elements)
size_t	cl::sycl::ext::oneapi::detail::reduGetMaxWGSize (std::shared_ptr< queue_impl > Queue, size_t LocalMemBytesPerWorkItem)
size_t	cl::sycl::ext::oneapi::detail::reduComputeWGSize (size_t NWorkItems, size_t MaxWGSize, size_t &NWorkGroups)
template<typename KernelFunc , int Dims, typename ReducerT >
void	cl::sycl::ext::oneapi::detail::reductionLoop (const range< Dims > &Range, ReducerT &Reducer, const nd_item< 1 > &NdId, KernelFunc &F)
	Called in device code. More...
template<typename KernelName , typename KernelType , int Dims, class Reduction >
std::enable_if_t< Reduction::has_fast_atomics >	cl::sycl::ext::oneapi::detail::reduCGFuncImpl (handler &CGH, KernelType KernelFunc, const range< Dims > &Range, const nd_range< 1 > &NDRange, Reduction &Redu)
template<typename KernelName , typename KernelType , int Dims, class Reduction >
std::enable_if_t<!Reduction::has_fast_atomics &&Reduction::has_fast_reduce >	cl::sycl::ext::oneapi::detail::reduCGFuncImpl (handler &CGH, KernelType KernelFunc, const range< Dims > &Range, const nd_range< 1 > &NDRange, Reduction &Redu)
template<typename KernelName , typename KernelType , int Dims, class Reduction >
std::enable_if_t<!Reduction::has_fast_atomics &&!Reduction::has_fast_reduce >	cl::sycl::ext::oneapi::detail::reduCGFuncImpl (handler &CGH, KernelType KernelFunc, const range< Dims > &Range, const nd_range< 1 > &NDRange, Reduction &Redu)
template<typename KernelName , typename KernelType , int Dims, class Reduction >
void	cl::sycl::ext::oneapi::detail::reduCGFunc (handler &CGH, KernelType KernelFunc, const range< Dims > &Range, size_t MaxWGSize, uint32_t NumConcurrentWorkGroups, Reduction &Redu)
template<typename KernelName , typename KernelType , int Dims, class Reduction , bool IsPow2WG>
enable_if_t< Reduction::has_fast_reduce &&Reduction::has_fast_atomics >	cl::sycl::ext::oneapi::detail::reduCGFuncImpl (handler &CGH, KernelType KernelFunc, const nd_range< Dims > &Range, Reduction &, typename Reduction::rw_accessor_type Out)
	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. More...
template<typename KernelName , typename KernelType , int Dims, class Reduction , bool IsPow2WG>
enable_if_t<!Reduction::has_fast_reduce &&Reduction::has_fast_atomics >	cl::sycl::ext::oneapi::detail::reduCGFuncImpl (handler &CGH, KernelType KernelFunc, const nd_range< Dims > &Range, Reduction &, typename Reduction::rw_accessor_type Out)
	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. More...
template<typename KernelName , typename KernelType , int Dims, class Reduction >
enable_if_t< Reduction::has_fast_atomics >	cl::sycl::ext::oneapi::detail::reduCGFunc (handler &CGH, KernelType KernelFunc, const nd_range< Dims > &Range, Reduction &Redu)
template<typename KernelName , typename KernelType , int Dims, class Reduction , bool IsPow2WG>
enable_if_t< Reduction::has_fast_reduce &&!Reduction::has_fast_atomics >	cl::sycl::ext::oneapi::detail::reduCGFuncImpl (handler &CGH, KernelType KernelFunc, const nd_range< Dims > &Range, Reduction &Redu, typename Reduction::rw_accessor_type Out)
	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. More...
template<typename KernelName , typename KernelType , int Dims, class Reduction , bool IsPow2WG>
enable_if_t<!Reduction::has_fast_reduce &&!Reduction::has_fast_atomics >	cl::sycl::ext::oneapi::detail::reduCGFuncImpl (handler &CGH, KernelType KernelFunc, const nd_range< Dims > &Range, Reduction &Redu, typename Reduction::rw_accessor_type Out)
	Implements a command group function that enqueues a kernel that calls user's lambda function. More...
template<typename KernelName , typename KernelType , int Dims, class Reduction >
enable_if_t<!Reduction::has_fast_atomics >	cl::sycl::ext::oneapi::detail::reduCGFunc (handler &CGH, KernelType KernelFunc, const nd_range< Dims > &Range, Reduction &Redu)
template<typename KernelName , typename KernelType , bool UniformWG, class Reduction , typename InputT , typename OutputT >
enable_if_t< Reduction::has_fast_reduce &&!Reduction::has_fast_atomics >	cl::sycl::ext::oneapi::detail::reduAuxCGFuncImpl (handler &CGH, size_t NWorkItems, size_t NWorkGroups, size_t WGSize, Reduction &Redu, InputT In, OutputT Out)
	Implements a command group function that enqueues a kernel that does one iteration of reduction of elements in each of work-groups. More...
template<typename KernelName , typename KernelType , bool UniformPow2WG, class Reduction , typename InputT , typename OutputT >
enable_if_t<!Reduction::has_fast_reduce &&!Reduction::has_fast_atomics >	cl::sycl::ext::oneapi::detail::reduAuxCGFuncImpl (handler &CGH, size_t NWorkItems, size_t NWorkGroups, size_t WGSize, Reduction &Redu, InputT In, OutputT Out)
	Implements a command group function that enqueues a kernel that does one iteration of reduction of elements in each of work-groups. More...
template<typename KernelName , typename KernelType , class Reduction >
enable_if_t<!Reduction::has_fast_atomics, size_t >	cl::sycl::ext::oneapi::detail::reduAuxCGFunc (handler &CGH, size_t NWorkItems, size_t MaxWGSize, Reduction &Redu)
	Implements a command group function that enqueues a kernel that does one iteration of reduction of elements in each of work-groups. More...
template<typename KernelName , class Reduction >
std::enable_if_t<!Reduction::is_usm >	cl::sycl::ext::oneapi::detail::reduSaveFinalResultToUserMem (handler &CGH, Reduction &Redu)
	Copies the final reduction result kept in read-write accessor to user's accessor. More...
template<typename KernelName , class Reduction >
std::enable_if_t< Reduction::is_usm >	cl::sycl::ext::oneapi::detail::reduSaveFinalResultToUserMem (handler &CGH, Reduction &Redu)
	Copies the final reduction result kept in read-write accessor to user's USM memory. More...
template<typename... Reductions, size_t... Is>
auto	cl::sycl::ext::oneapi::detail::createReduLocalAccs (size_t Size, handler &CGH, std::index_sequence< Is... >)
	For the given 'Reductions' types pack and indices enumerating only the reductions for which a local accessors are needed, this function creates those local accessors and returns a tuple consisting of them. More...
template<bool IsOneWG, typename... Reductions, size_t... Is>
auto	cl::sycl::ext::oneapi::detail::createReduOutAccs (size_t NWorkGroups, handler &CGH, std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... >)
	For the given 'Reductions' types pack and indices enumerating them this function either creates new temporary accessors for partial sums (if IsOneWG is false) or returns user's accessor/USM-pointer if (IsOneWG is true). More...
template<typename... Reductions, size_t... Is>
auto	cl::sycl::ext::oneapi::detail::getReadAccsToPreviousPartialReds (handler &CGH, std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... >)
	For the given 'Reductions' types pack and indices enumerating them this function returns accessors to buffers holding partial sums generated in the previous kernel invocation. More...
template<typename... Reductions, size_t... Is>
ReduTupleT< typename Reductions::result_type... >	cl::sycl::ext::oneapi::detail::getReduIdentities (std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... >)
template<typename... Reductions, size_t... Is>
ReduTupleT< typename Reductions::binary_operation... >	cl::sycl::ext::oneapi::detail::getReduBOPs (std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... >)
template<typename... Reductions, size_t... Is>
std::array< bool, sizeof...(Reductions)>	cl::sycl::ext::oneapi::detail::getInitToIdentityProperties (std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... >)
template<typename... Reductions, size_t... Is>
std::tuple< typename Reductions::reducer_type... >	cl::sycl::ext::oneapi::detail::createReducers (ReduTupleT< typename Reductions::result_type... > Identities, ReduTupleT< typename Reductions::binary_operation... > BOPsTuple, std::index_sequence< Is... >)
template<typename KernelType , int Dims, typename... ReducerT, size_t... Is>
void	cl::sycl::ext::oneapi::detail::callReduUserKernelFunc (KernelType KernelFunc, nd_item< Dims > NDIt, std::tuple< ReducerT... > &Reducers, std::index_sequence< Is... >)
template<bool Pow2WG, typename... LocalAccT, typename... ReducerT, typename... ResultT, size_t... Is>
void	cl::sycl::ext::oneapi::detail::initReduLocalAccs (size_t LID, size_t WGSize, ReduTupleT< LocalAccT... > LocalAccs, const std::tuple< ReducerT... > &Reducers, ReduTupleT< ResultT... > Identities, std::index_sequence< Is... >)
template<bool UniformPow2WG, typename... LocalAccT, typename... InputAccT, typename... ResultT, size_t... Is>
void	cl::sycl::ext::oneapi::detail::initReduLocalAccs (size_t LID, size_t GID, size_t NWorkItems, size_t WGSize, ReduTupleT< InputAccT... > LocalAccs, ReduTupleT< LocalAccT... > InputAccs, ReduTupleT< ResultT... > Identities, std::index_sequence< Is... >)
template<typename... LocalAccT, typename... BOPsT, size_t... Is>
void	cl::sycl::ext::oneapi::detail::reduceReduLocalAccs (size_t IndexA, size_t IndexB, ReduTupleT< LocalAccT... > LocalAccs, ReduTupleT< BOPsT... > BOPs, std::index_sequence< Is... >)
template<bool Pow2WG, bool IsOneWG, typename... Reductions, typename... OutAccT, typename... LocalAccT, typename... BOPsT, typename... Ts, size_t... Is>
void	cl::sycl::ext::oneapi::detail::writeReduSumsToOutAccs (size_t OutAccIndex, size_t WGSize, std::tuple< Reductions... > *, ReduTupleT< OutAccT... > OutAccs, ReduTupleT< LocalAccT... > LocalAccs, ReduTupleT< BOPsT... > BOPs, ReduTupleT< Ts... > IdentityVals, std::array< bool, sizeof...(Reductions)> IsInitializeToIdentity, std::index_sequence< Is... >)
constexpr std::index_sequence	cl::sycl::ext::oneapi::detail::concat_sequences (std::index_sequence<>)
template<size_t I>
constexpr std::index_sequence< I >	cl::sycl::ext::oneapi::detail::concat_sequences (std::index_sequence< I >)
template<size_t... Is, size_t... Js>
constexpr std::index_sequence< Is..., Js... >	cl::sycl::ext::oneapi::detail::concat_sequences (std::index_sequence< Is... >, std::index_sequence< Js... >)
template<size_t... Is, size_t... Js, class... Rs>
constexpr auto	cl::sycl::ext::oneapi::detail::concat_sequences (std::index_sequence< Is... >, std::index_sequence< Js... >, Rs...)
template<typename... T, typename FunctorT , size_t... Is, std::enable_if_t<(sizeof...(Is) > 0), int > Z = 0>
constexpr auto	cl::sycl::ext::oneapi::detail::filterSequenceHelper (FunctorT, std::index_sequence< Is... >)
	For each index 'I' from the given indices pack 'Is' this function initially creates a number of short index_sequences, where each of such short index sequences is either empty (if the given Functor returns false for the type T[I]) or 1 element 'I' (otherwise). More...
template<typename... T, typename FunctorT , size_t... Is>
constexpr auto	cl::sycl::ext::oneapi::detail::filterSequence (FunctorT F, std::index_sequence< Is... > Indices)
	For each index 'I' from the given indices pack 'Is' this function returns an index sequence consisting of only those 'I's for which the 'FunctorT' applied to 'T[I]' returns true. More...
template<bool Pow2WG, bool IsOneWG, typename... Reductions, int Dims, typename... LocalAccT, typename... OutAccT, typename... ReducerT, typename... Ts, typename... BOPsT, size_t... Is>
void	cl::sycl::ext::oneapi::detail::reduCGFuncImplScalar (nd_item< Dims > NDIt, ReduTupleT< LocalAccT... > LocalAccsTuple, ReduTupleT< OutAccT... > OutAccsTuple, std::tuple< ReducerT... > &ReducersTuple, ReduTupleT< Ts... > IdentitiesTuple, ReduTupleT< BOPsT... > BOPsTuple, std::array< bool, sizeof...(Reductions)> InitToIdentityProps, std::index_sequence< Is... > ReduIndices)
	All scalar reductions are processed together; there is one loop of log2(N) steps, and each reduction uses its own storage. More...
template<bool Pow2WG, bool IsOneWG, typename Reduction , int Dims, typename LocalAccT , typename OutAccT , typename ReducerT , typename T , typename BOPT >
void	cl::sycl::ext::oneapi::detail::reduCGFuncImplArrayHelper (nd_item< Dims > NDIt, LocalAccT LocalReds, OutAccT Out, ReducerT &Reducer, T Identity, BOPT BOp, bool IsInitializeToIdentity)
	Each array reduction is processed separately. More...
template<bool Pow2WG, bool IsOneWG, typename... Reductions, int Dims, typename... LocalAccT, typename... OutAccT, typename... ReducerT, typename... Ts, typename... BOPsT, size_t... Is>
void	cl::sycl::ext::oneapi::detail::reduCGFuncImplArray (nd_item< Dims > NDIt, ReduTupleT< LocalAccT... > LocalAccsTuple, ReduTupleT< OutAccT... > OutAccsTuple, std::tuple< ReducerT... > &ReducersTuple, ReduTupleT< Ts... > IdentitiesTuple, ReduTupleT< BOPsT... > BOPsTuple, std::array< bool, sizeof...(Reductions)> InitToIdentityProps, std::index_sequence< Is... >)
template<typename KernelName , bool Pow2WG, bool IsOneWG, typename KernelType , int Dims, typename... Reductions, size_t... Is>
void	cl::sycl::ext::oneapi::detail::reduCGFuncImpl (handler &CGH, KernelType KernelFunc, const nd_range< Dims > &Range, std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... > ReduIndices)
template<typename KernelName , typename KernelType , int Dims, typename... Reductions, size_t... Is>
void	cl::sycl::ext::oneapi::detail::reduCGFunc (handler &CGH, KernelType KernelFunc, const nd_range< Dims > &Range, std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... >)
template<typename KernelName , typename KernelType , int Dims, class Reduction >
std::enable_if_t< Reduction::has_atomic_add_float64 >	cl::sycl::ext::oneapi::detail::reduCGFuncImplAtomic64 (handler &CGH, KernelType KernelFunc, const nd_range< Dims > &Range, Reduction &, typename Reduction::rw_accessor_type Out)
template<typename KernelName , typename KernelType , int Dims, class Reduction >
enable_if_t< Reduction::has_atomic_add_float64 >	cl::sycl::ext::oneapi::detail::reduCGFuncAtomic64 (handler &CGH, KernelType KernelFunc, const nd_range< Dims > &Range, Reduction &Redu)
void	cl::sycl::ext::oneapi::detail::associateReduAccsWithHandlerHelper (handler &)
template<typename ReductionT >
void	cl::sycl::ext::oneapi::detail::associateReduAccsWithHandlerHelper (handler &CGH, ReductionT &Redu)
template<typename ReductionT , typename... RestT, enable_if_t<(sizeof...(RestT) > 0), int > Z = 0>
void	cl::sycl::ext::oneapi::detail::associateReduAccsWithHandlerHelper (handler &CGH, ReductionT &Redu, RestT &... Rest)
template<typename... Reductions, size_t... Is>
void	cl::sycl::ext::oneapi::detail::associateReduAccsWithHandler (handler &CGH, std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... >)
template<bool UniformPow2WG, bool IsOneWG, typename... Reductions, int Dims, typename... LocalAccT, typename... InAccT, typename... OutAccT, typename... Ts, typename... BOPsT, size_t... Is>
void	cl::sycl::ext::oneapi::detail::reduAuxCGFuncImplScalar (nd_item< Dims > NDIt, size_t LID, size_t GID, size_t NWorkItems, size_t WGSize, ReduTupleT< LocalAccT... > LocalAccsTuple, ReduTupleT< InAccT... > InAccsTuple, ReduTupleT< OutAccT... > OutAccsTuple, ReduTupleT< Ts... > IdentitiesTuple, ReduTupleT< BOPsT... > BOPsTuple, std::array< bool, sizeof...(Reductions)> InitToIdentityProps, std::index_sequence< Is... > ReduIndices)
	All scalar reductions are processed together; there is one loop of log2(N) steps, and each reduction uses its own storage. More...
template<bool UniformPow2WG, bool IsOneWG, typename Reduction , int Dims, typename LocalAccT , typename InAccT , typename OutAccT , typename T , typename BOPT >
void	cl::sycl::ext::oneapi::detail::reduAuxCGFuncImplArrayHelper (nd_item< Dims > NDIt, size_t LID, size_t GID, size_t NWorkItems, size_t WGSize, LocalAccT LocalReds, InAccT In, OutAccT Out, T Identity, BOPT BOp, bool IsInitializeToIdentity)
template<bool UniformPow2WG, bool IsOneWG, typename... Reductions, int Dims, typename... LocalAccT, typename... InAccT, typename... OutAccT, typename... Ts, typename... BOPsT, size_t... Is>
void	cl::sycl::ext::oneapi::detail::reduAuxCGFuncImplArray (nd_item< Dims > NDIt, size_t LID, size_t GID, size_t NWorkItems, size_t WGSize, ReduTupleT< LocalAccT... > LocalAccsTuple, ReduTupleT< InAccT... > InAccsTuple, ReduTupleT< OutAccT... > OutAccsTuple, ReduTupleT< Ts... > IdentitiesTuple, ReduTupleT< BOPsT... > BOPsTuple, std::array< bool, sizeof...(Reductions)> InitToIdentityProps, std::index_sequence< Is... >)
template<typename KernelName , typename KernelType , bool UniformPow2WG, bool IsOneWG, typename... Reductions, size_t... Is>
void	cl::sycl::ext::oneapi::detail::reduAuxCGFuncImpl (handler &CGH, size_t NWorkItems, size_t NWorkGroups, size_t WGSize, std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... > ReduIndices)
template<typename KernelName , typename KernelType , typename... Reductions, size_t... Is>
size_t	cl::sycl::ext::oneapi::detail::reduAuxCGFunc (handler &CGH, size_t NWorkItems, size_t MaxWGSize, std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... >)
void	cl::sycl::ext::oneapi::detail::reduSaveFinalResultToUserMemHelper (std::vector< event > &, std::shared_ptr< detail::queue_impl >, bool)
template<typename Reduction , typename... RestT>
std::enable_if_t< Reduction::is_usm >	cl::sycl::ext::oneapi::detail::reduSaveFinalResultToUserMemHelper (std::vector< event > &Events, std::shared_ptr< detail::queue_impl > Queue, bool IsHost, Reduction &, RestT... Rest)
template<typename Reduction , typename... RestT>
std::enable_if_t<!Reduction::is_usm >	cl::sycl::ext::oneapi::detail::reduSaveFinalResultToUserMemHelper (std::vector< event > &Events, std::shared_ptr< detail::queue_impl > Queue, bool IsHost, Reduction &Redu, RestT... Rest)
template<typename... Reduction, size_t... Is>
std::shared_ptr< event >	cl::sycl::ext::oneapi::detail::reduSaveFinalResultToUserMem (std::shared_ptr< detail::queue_impl > Queue, bool IsHost, std::tuple< Reduction... > &ReduTuple, std::index_sequence< Is... >)
	Creates additional kernels that copy the accumulated/final results from reductions accessors to either user's accessor or user's USM memory. More...
template<typename Reduction >
size_t	cl::sycl::ext::oneapi::detail::reduGetMemPerWorkItemHelper (Reduction &)
template<typename Reduction , typename... RestT>
size_t	cl::sycl::ext::oneapi::detail::reduGetMemPerWorkItemHelper (Reduction &, RestT... Rest)
template<typename... ReductionT, size_t... Is>
size_t	cl::sycl::ext::oneapi::detail::reduGetMemPerWorkItem (std::tuple< ReductionT... > &ReduTuple, std::index_sequence< Is... >)
template<typename TupleT , std::size_t... Is>
std::tuple< std::tuple_element_t< Is, TupleT >... >	cl::sycl::ext::oneapi::detail::tuple_select_elements (TupleT Tuple, std::index_sequence< Is... >)
	Utility function: for the given tuple. More...
template<typename T , class BinaryOperation , int Dims, access::mode AccMode, access::placeholder IsPH>
detail::reduction_impl< T, BinaryOperation, 0, 1, detail::default_reduction_algorithm< false, IsPH, Dims > >	cl::sycl::ext::oneapi::reduction (accessor< T, Dims, AccMode, access::target::device, IsPH > &Acc, const T &Identity, BinaryOperation BOp)
	Creates and returns an object implementing the reduction functionality. More...
template<typename T , class BinaryOperation , int Dims, access::mode AccMode, access::placeholder IsPH>
std::enable_if_t< detail::IsKnownIdentityOp< T, BinaryOperation >::value, detail::reduction_impl< T, BinaryOperation, 0, 1, detail::default_reduction_algorithm< false, IsPH, Dims > > >	cl::sycl::ext::oneapi::reduction (accessor< T, Dims, AccMode, access::target::device, IsPH > &Acc, BinaryOperation)
	Creates and returns an object implementing the reduction functionality. More...
template<typename T , class BinaryOperation >
detail::reduction_impl< T, BinaryOperation, 0, 1, detail::default_reduction_algorithm< true, access::placeholder::false_t, 1 > >	cl::sycl::ext::oneapi::reduction (T *VarPtr, const T &Identity, BinaryOperation BOp)
	Creates and returns an object implementing the reduction functionality. More...
template<typename T , class BinaryOperation >
std::enable_if_t< detail::IsKnownIdentityOp< T, BinaryOperation >::value, detail::reduction_impl< T, BinaryOperation, 0, 1, detail::default_reduction_algorithm< true, access::placeholder::false_t, 1 > > >	cl::sycl::ext::oneapi::reduction (T *VarPtr, BinaryOperation)
	Creates and returns an object implementing the reduction functionality. More...

Variables
template<typename BinaryOperation , typename AccumulatorT >
__SYCL_INLINE_CONSTEXPR bool	cl::sycl::ext::oneapi::has_known_identity_v
template<typename BinaryOperation , typename AccumulatorT >
__SYCL_INLINE_CONSTEXPR AccumulatorT	cl::sycl::ext::oneapi::known_identity_v