DPC++ Runtime
Runtime libraries for oneAPI DPC++
reduction.hpp File Reference
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  sycl::_V1::detail::reduction_impl_base
 Base non-template class which is a base class for all reduction implementation classes. More...
 
struct  sycl::_V1::detail::IsReduction< T >
 Predicate returning true if a type is a reduction. More...
 
struct  sycl::_V1::detail::AreAllButLastReductions< FirstT, RestT >
 Predicate returning true if all template type parameters except the last one are reductions. More...
 
struct  sycl::_V1::detail::AreAllButLastReductions< T >
 Helper specialization of AreAllButLastReductions for one element only. More...
 
struct  sycl::_V1::detail::ReducerTraits< reducer< T, BinaryOperation, Dims, Extent, View, Subst > >
 
class  sycl::_V1::detail::combiner< Reducer >
 Use CRTP to avoid redefining shorthand operators in terms of combine. More...
 
class  sycl::_V1::reducer< T, BinaryOperation, Dims, Extent, View, std::enable_if_t< Dims==0 &&Extent==1 &&View==false &&!detail::IsKnownIdentityOp< T, BinaryOperation >::value > >
 Specialization of the generic class 'reducer'. More...
 
class  sycl::_V1::reducer< T, BinaryOperation, Dims, Extent, View, std::enable_if_t< Dims==0 &&Extent==1 &&View==false &&detail::IsKnownIdentityOp< T, BinaryOperation >::value > >
 Specialization of the generic class 'reducer'. More...
 
class  sycl::_V1::reducer< T, BinaryOperation, Dims, Extent, View, std::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  sycl::_V1::reducer< T, BinaryOperation, Dims, Extent, View, std::enable_if_t< Dims==1 &&View==false &&!detail::IsKnownIdentityOp< T, BinaryOperation >::value > >
 Specialization of 'reducer' class for array reductions exposing the subscript operator. More...
 
class  sycl::_V1::reducer< T, BinaryOperation, Dims, Extent, View, std::enable_if_t< Dims==1 &&View==false &&detail::IsKnownIdentityOp< T, BinaryOperation >::value > >
 Specialization of 'reducer' class for array reductions accepting a span in cases where the identity value is known. More...
 
class  sycl::_V1::detail::reduction_impl_common< T, BinaryOperation >
 Templated class for common functionality of all reduction implementation classes. More...
 
struct  sycl::_V1::detail::data_dim_t< T >
 
struct  sycl::_V1::detail::data_dim_t< accessor< T, AccessorDims, Mode, access::target::device, IsPH, PropList > >
 
struct  sycl::_V1::detail::get_red_t< T * >
 
struct  sycl::_V1::detail::get_red_t< accessor< T, AccessorDims, Mode, access::target::device, IsPH, PropList > >
 
class  sycl::_V1::detail::reduction_impl_algo< T, BinaryOperation, Dims, Extent, RedOutVar >
 
class  sycl::_V1::detail::reduction_impl< T, BinaryOperation, Dims, Extent, RedOutVar >
 This class encapsulates the reduction variable/accessor, the reduction operator and an optional operator identity. More...
 
struct  sycl::_V1::detail::NDRangeReduction< reduction::strategy::local_atomic_and_atomic_cross_wg >
 
struct  sycl::_V1::detail::NDRangeReduction< reduction::strategy::group_reduce_and_last_wg_detection >
 
struct  sycl::_V1::detail::NDRangeReduction< reduction::strategy::range_basic >
 
struct  sycl::_V1::detail::NDRangeReduction< reduction::strategy::group_reduce_and_atomic_cross_wg >
 
struct  sycl::_V1::detail::NDRangeReduction< reduction::strategy::local_mem_tree_and_atomic_cross_wg >
 
struct  sycl::_V1::detail::NDRangeReduction< reduction::strategy::group_reduce_and_multiple_kernels >
 
struct  sycl::_V1::detail::NDRangeReduction< reduction::strategy::basic >
 
struct  sycl::_V1::detail::IsNonUsmReductionPredicate
 
struct  sycl::_V1::detail::IsNonUsmReductionPredicate::Func< T >
 
struct  sycl::_V1::detail::EmptyReductionPredicate
 
struct  sycl::_V1::detail::EmptyReductionPredicate::Func< T >
 
struct  sycl::_V1::detail::FilterElement< Cond, I >
 
struct  sycl::_V1::detail::IsScalarReduction
 
struct  sycl::_V1::detail::IsScalarReduction::Func< Reduction >
 
struct  sycl::_V1::detail::IsArrayReduction
 
struct  sycl::_V1::detail::IsArrayReduction::Func< Reduction >
 
struct  sycl::_V1::detail::NDRangeReduction< reduction::strategy::multi >
 
struct  sycl::_V1::detail::NDRangeReduction< reduction::strategy::auto_select >
 

Namespaces

 sycl
 ---— Error handling, matching OpenCL plugin semantics.
 
 sycl::_V1
 
 sycl::_V1::detail
 
 sycl::_V1::detail::reduction
 
 sycl::_V1::detail::reduction::main_krn
 
 sycl::_V1::detail::reduction::aux_krn
 

Typedefs

template<typename T , class BinaryOperation >
using sycl::_V1::detail::IsReduOptForFastAtomicFetch = std::bool_constant<((is_sgenfloat< T >::value &&sizeof(T)==4)||is_sgeninteger< T >::value) &&IsValidAtomicType< T >::value &&(IsPlus< T, BinaryOperation >::value||IsMinimum< T, BinaryOperation >::value||IsMaximum< T, BinaryOperation >::value||IsBitOR< T, BinaryOperation >::value||IsBitXOR< T, BinaryOperation >::value||IsBitAND< T, BinaryOperation >::value)>
 
template<typename T , class BinaryOperation >
using sycl::_V1::detail::IsReduOptForAtomic64Op = std::bool_constant<(IsPlus< T, BinaryOperation >::value||IsMinimum< T, BinaryOperation >::value||IsMaximum< T, BinaryOperation >::value) &&is_sgenfloat< T >::value &&sizeof(T)==8 >
 
template<typename T , class BinaryOperation >
using sycl::_V1::detail::IsReduOptForFastReduce = std::bool_constant<((is_sgeninteger< T >::value &&(sizeof(T)==4||sizeof(T)==8))||is_sgenfloat< T >::value) &&(IsPlus< T, BinaryOperation >::value||IsMinimum< T, BinaryOperation >::value||IsMaximum< T, BinaryOperation >::value)>
 
template<typename... Ts>
using sycl::_V1::detail::ReduTupleT = sycl::detail::tuple< Ts... >
 
template<template< typename, reduction::strategy, typename... > class MainOrAux, class KernelName , reduction::strategy Strategy, class... Ts>
using sycl::_V1::detail::__sycl_reduction_kernel = std::conditional_t< std::is_same< KernelName, auto_name >::value, auto_name, MainOrAux< KernelName, Strategy, Ts... > >
 A helper to pass undefined (sycl::detail::auto_name) names unmodified. More...
 

Functions

template<typename... Ts>
ReduTupleT< Ts... > sycl::_V1::detail::makeReduTupleT (Ts... Elements)
 
size_t sycl::_V1::detail::reduGetMaxWGSize (std::shared_ptr< queue_impl > Queue, size_t LocalMemBytesPerWorkItem)
 
size_t sycl::_V1::detail::reduComputeWGSize (size_t NWorkItems, size_t MaxWGSize, size_t &NWorkGroups)
 
size_t sycl::_V1::detail::reduGetPreferredWGSize (std::shared_ptr< queue_impl > &Queue, size_t LocalMemBytesPerWorkItem)
 
template<class BinaryOp , int Dims, size_t Extent, typename RedOutVar , typename... RestTy>
auto sycl::_V1::detail::make_reduction (RedOutVar RedVar, RestTy &&...Rest)
 
void sycl::_V1::detail::reduction::finalizeHandler (handler &CGH)
 
template<class FunctorTy >
void sycl::_V1::detail::reduction::withAuxHandler (handler &CGH, FunctorTy Func)
 
template<typename KernelName , class Reduction >
void sycl::_V1::detail::reduSaveFinalResultToUserMem (handler &CGH, Reduction &Redu)
 Copies the final reduction result kept in read-write accessor to user's USM memory. More...
 
template<bool IsOneWG, typename... Reductions, size_t... Is>
auto sycl::_V1::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... LocalAccT, typename... BOPsT, size_t... Is>
void sycl::_V1::detail::reduceReduLocalAccs (size_t IndexA, size_t IndexB, ReduTupleT< LocalAccT... > LocalAccs, ReduTupleT< BOPsT... > BOPs, std::index_sequence< Is... >)
 
template<typename... Reductions, typename... OutAccT, typename... LocalAccT, typename... BOPsT, typename... Ts, size_t... Is>
void sycl::_V1::detail::writeReduSumsToOutAccs (bool Pow2WG, bool IsOneWG, size_t OutAccIndex, size_t WGSize, 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 sycl::_V1::detail::concat_sequences (std::index_sequence<>)
 
template<size_t I>
constexpr std::index_sequence< I > sycl::_V1::detail::concat_sequences (std::index_sequence< I >)
 
template<size_t... Is, size_t... Js>
constexpr std::index_sequence< Is..., Js... > sycl::_V1::detail::concat_sequences (std::index_sequence< Is... >, std::index_sequence< Js... >)
 
template<size_t... Is, size_t... Js, class... Rs>
constexpr auto sycl::_V1::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 sycl::_V1::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 sycl::_V1::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... Reductions, int Dims, typename... LocalAccT, typename... OutAccT, typename... ReducerT, typename... Ts, typename... BOPsT, size_t... Is>
void sycl::_V1::detail::reduCGFuncImplScalar (bool Pow2WG, bool IsOneWG, 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<typename Reduction , int Dims, typename LocalAccT , typename OutAccT , typename ReducerT , typename T , typename BOPT >
void sycl::_V1::detail::reduCGFuncImplArrayHelper (bool Pow2WG, bool IsOneWG, nd_item< Dims > NDIt, LocalAccT LocalReds, OutAccT Out, ReducerT &Reducer, T Identity, BOPT BOp, bool IsInitializeToIdentity)
 Each array reduction is processed separately. More...
 
template<typename... Reductions, int Dims, typename... LocalAccT, typename... OutAccT, typename... ReducerT, typename... Ts, typename... BOPsT, size_t... Is>
void sycl::_V1::detail::reduCGFuncImplArray (bool Pow2WG, bool IsOneWG, 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 , typename KernelType , int Dims, typename PropertiesT , typename... Reductions, size_t... Is>
void sycl::_V1::detail::reduCGFuncMulti (handler &CGH, KernelType KernelFunc, const nd_range< Dims > &Range, PropertiesT Properties, std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... > ReduIndices)
 
template<typename... Reductions, size_t... Is>
void sycl::_V1::detail::associateReduAccsWithHandler (handler &CGH, std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... >)
 
template<typename... Reductions, int Dims, typename... LocalAccT, typename... InAccT, typename... OutAccT, typename... Ts, typename... BOPsT, size_t... Is>
void sycl::_V1::detail::reduAuxCGFuncImplScalar (bool UniformPow2WG, bool IsOneWG, 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<typename Reduction , int Dims, typename LocalAccT , typename InAccT , typename OutAccT , typename T , typename BOPT >
void sycl::_V1::detail::reduAuxCGFuncImplArrayHelper (bool UniformPow2WG, bool IsOneWG, 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<typename... Reductions, int Dims, typename... LocalAccT, typename... InAccT, typename... OutAccT, typename... Ts, typename... BOPsT, size_t... Is>
void sycl::_V1::detail::reduAuxCGFuncImplArray (bool UniformPow2WG, bool IsOneWG, 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 , typename... Reductions, size_t... Is>
size_t sycl::_V1::detail::reduAuxCGFunc (handler &CGH, size_t NWorkItems, size_t MaxWGSize, std::tuple< Reductions... > &ReduTuple, std::index_sequence< Is... > ReduIndices)
 
template<typename Reduction >
size_t sycl::_V1::detail::reduGetMemPerWorkItemHelper (Reduction &)
 
template<typename Reduction , typename... RestT>
size_t sycl::_V1::detail::reduGetMemPerWorkItemHelper (Reduction &, RestT... Rest)
 
template<typename... ReductionT, size_t... Is>
size_t sycl::_V1::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 >... > sycl::_V1::detail::tuple_select_elements (TupleT Tuple, std::index_sequence< Is... >)
 Utility function: for the given tuple. More...
 
template<typename KernelName , reduction::strategy Strategy, int Dims, typename PropertiesT , typename... RestT>
void sycl::_V1::detail::reduction_parallel_for (handler &CGH, nd_range< Dims > NDRange, PropertiesT Properties, RestT... Rest)
 
uint32_t sycl::_V1::detail::reduGetMaxNumConcurrentWorkGroups (std::shared_ptr< queue_impl > Queue)
 
template<typename KernelName , reduction::strategy Strategy, int Dims, typename PropertiesT , typename... RestT>
void sycl::_V1::detail::reduction_parallel_for (handler &CGH, range< Dims > Range, PropertiesT Properties, RestT... Rest)
 
template<typename T , typename AllocatorT , typename BinaryOperation , typename = std::enable_if_t<has_known_identity<BinaryOperation, T>::value>>
auto sycl::_V1::reduction (buffer< T, 1, AllocatorT > Var, handler &CGH, BinaryOperation, const property_list &PropList={})
 Constructs a reduction object using the given buffer Var, handler CGH, reduction operation Combiner, and optional reduction properties. More...
 
template<typename T , typename AllocatorT , typename BinaryOperation , typename = std::enable_if_t<!has_known_identity<BinaryOperation, T>::value>>
detail::reduction_impl< T, BinaryOperation, 0, 1, accessor< T, 1, access::mode::read_write, access::target::device, access::placeholder::true_t, ext::oneapi::accessor_property_list<> > > sycl::_V1::reduction (buffer< T, 1, AllocatorT >, handler &, BinaryOperation, const property_list &PropList={})
 Constructs a reduction object using the given buffer Var, handler CGH, reduction operation Combiner, and optional reduction properties. More...
 
template<typename T , typename BinaryOperation , typename = std::enable_if_t<has_known_identity<BinaryOperation, T>::value>>
auto sycl::_V1::reduction (T *Var, BinaryOperation, const property_list &PropList={})
 Constructs a reduction object using the reduction variable referenced by the given USM pointer Var, handler CGH, reduction operation Combiner, and optional reduction properties. More...
 
template<typename T , typename BinaryOperation , typename = std::enable_if_t<!has_known_identity<BinaryOperation, T>::value>>
detail::reduction_impl< T, BinaryOperation, 0, 1, T * > sycl::_V1::reduction (T *, BinaryOperation, const property_list &PropList={})
 Constructs a reduction object using the reduction variable referenced by the given USM pointer Var, handler CGH, reduction operation Combiner, and optional reduction properties. More...
 
template<typename T , typename AllocatorT , typename BinaryOperation >
auto sycl::_V1::reduction (buffer< T, 1, AllocatorT > Var, handler &CGH, const T &Identity, BinaryOperation Combiner, const property_list &PropList={})
 Constructs a reduction object using the given buffer Var, handler CGH, reduction identity value Identity, reduction operation Combiner, and optional reduction properties. More...
 
template<typename T , typename BinaryOperation >
auto sycl::_V1::reduction (T *Var, const T &Identity, BinaryOperation Combiner, const property_list &PropList={})
 Constructs a reduction object using the reduction variable referenced by the given USM pointer Var, reduction identity value Identity, binary operation Combiner, and optional reduction properties. More...
 
template<typename T , size_t Extent, typename BinaryOperation , typename = std::enable_if_t<Extent != dynamic_extent && has_known_identity<BinaryOperation, T>::value>>
auto sycl::_V1::reduction (span< T, Extent > Span, BinaryOperation, const property_list &PropList={})
 Constructs a reduction object using the reduction variable referenced by the given sycl::span Span, reduction operation Combiner, and optional reduction properties. More...
 
template<typename T , size_t Extent, typename BinaryOperation , typename = std::enable_if_t<Extent != dynamic_extent && !has_known_identity<BinaryOperation, T>::value>>
detail::reduction_impl< T, BinaryOperation, 1, Extent, T * > sycl::_V1::reduction (span< T, Extent >, BinaryOperation, const property_list &PropList={})
 Constructs a reduction object using the reduction variable referenced by the given sycl::span Span, reduction operation Combiner, and optional reduction properties. More...
 
template<typename T , size_t Extent, typename BinaryOperation , typename = std::enable_if_t<Extent != dynamic_extent>>
auto sycl::_V1::reduction (span< T, Extent > Span, const T &Identity, BinaryOperation Combiner, const property_list &PropList={})
 Constructs a reduction object using the reduction variable referenced by the given sycl::span Span, reduction identity value Identity, reduction operation Combiner, and optional reduction properties. More...