#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, Subst >
	Class that is used to represent objects that are passed to user's lambda functions and representing users' reduction variable. More...

class	cl::sycl::ext::oneapi::detail::reducer< T, BinaryOperation, enable_if_t< IsKnownIdentityOp< T, BinaryOperation >::value > >
	Specialization of the generic class 'reducer'. 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...

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, IsUSM, IsPlaceholder >
	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::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<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<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, Dims, false, IsPH >	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, Dims, false, IsPH > >	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, 1, true >	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, 1, true > >	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

Classes

Namespaces

Typedefs

Functions

Variables