Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range > Class Template Reference

#include <TensorReductionSycl.h>

Public Types
typedef Evaluator::CoeffReturnType	CoeffReturnType
typedef Evaluator::Index	Index
typedef OpDefiner< OpType, typename Evaluator::CoeffReturnType, Index,(Evaluator::ReducerTraits::PacketAccess &Evaluator::InputPacketAccess)>	OpDef
typedef OpDef::type	Op
typedef Evaluator::EvaluatorPointerType	EvaluatorPointerType
typedef Evaluator::PacketReturnType	PacketReturnType
typedef std::conditional_t<(Evaluator::ReducerTraits::PacketAccess &Evaluator::InputPacketAccess), PacketReturnType, CoeffReturnType >	OutType
typedef cl::sycl::accessor< OutType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local >	LocalAccessor

Public Member Functions
	FullReductionKernelFunctor (LocalAccessor scratch_, Evaluator evaluator_, EvaluatorPointerType final_output_, Index rng_, OpType op_)
void	operator() (cl::sycl::nd_item< 1 > itemID) const
template<bool Vect = (Evaluator::ReducerTraits::PacketAccess & Evaluator::InputPacketAccess)>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t< Vect >	compute_reduction (const cl::sycl::nd_item< 1 > &itemID) const
template<bool Vect = (Evaluator::ReducerTraits::PacketAccess & Evaluator::InputPacketAccess)>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<!Vect >	compute_reduction (const cl::sycl::nd_item< 1 > &itemID) const

Public Attributes
LocalAccessor	scratch
Evaluator	evaluator
EvaluatorPointerType	final_output
Index	rng
Op	op

Member Typedef Documentation

CoeffReturnType

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

typedef Evaluator::CoeffReturnType Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::CoeffReturnType

EvaluatorPointerType

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

typedef Evaluator::EvaluatorPointerType Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::EvaluatorPointerType

Index

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

typedef Evaluator::Index Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::Index

LocalAccessor

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

typedef cl::sycl::accessor<OutType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::LocalAccessor

Op

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

typedef OpDef::type Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::Op

OpDef

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

typedef OpDefiner<OpType, typename Evaluator::CoeffReturnType, Index, (Evaluator::ReducerTraits::PacketAccess & Evaluator::InputPacketAccess)> Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::OpDef

OutType

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

typedef std::conditional_t<(Evaluator::ReducerTraits::PacketAccess & Evaluator::InputPacketAccess), PacketReturnType, CoeffReturnType> Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::OutType

PacketReturnType

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

typedef Evaluator::PacketReturnType Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::PacketReturnType

Constructor & Destructor Documentation

FullReductionKernelFunctor()

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::FullReductionKernelFunctor ( LocalAccessor  scratch_, Evaluator  evaluator_, EvaluatorPointerType  final_output_, Index  rng_, OpType  op_  )

inline

      : scratch(scratch_), evaluator(evaluator_), final_output(final_output_), rng(rng_), op(OpDef::get_op(op_)) {}

Member Function Documentation

compute_reduction() [1/2]

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

template<bool Vect = (Evaluator::ReducerTraits::PacketAccess & Evaluator::InputPacketAccess)>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<Vect> Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::compute_reduction ( const cl::sycl::nd_item< 1 > &  itemID ) const

inline

                                            {
    auto output_ptr = final_output;
    Index VectorizedRange = (rng / Evaluator::PacketSize) * Evaluator::PacketSize;
    Index globalid = itemID.get_global_id(0);
    Index localid = itemID.get_local_id(0);
    Index step = Evaluator::PacketSize * itemID.get_global_range(0);
    Index start = Evaluator::PacketSize * globalid;
    // vectorizable parts
    PacketReturnType packetAccumulator = op.template initializePacket<PacketReturnType>();
    for (Index i = start; i < VectorizedRange; i += step) {
      op.template reducePacket<PacketReturnType>(evaluator.impl().template packet<Unaligned>(i), &packetAccumulator);
    }
    globalid += VectorizedRange;
    // non vectorizable parts
    for (Index i = globalid; i < rng; i += itemID.get_global_range(0)) {
      op.template reducePacket<PacketReturnType>(
          ::Eigen::TensorSycl::internal::PacketWrapper<PacketReturnType, Evaluator::PacketSize>::convert_to_packet_type(
              evaluator.impl().coeff(i), op.initialize()),
          &packetAccumulator);
    }
    scratch[localid] = packetAccumulator =
        OpDef::finalise_op(op.template finalizePacket<PacketReturnType>(packetAccumulator), rng);
    // reduction parts // Local size is always power of 2
    EIGEN_UNROLL_LOOP
    for (Index offset = local_range / 2; offset > 0; offset /= 2) {
      itemID.barrier(cl::sycl::access::fence_space::local_space);
      if (localid < offset) {
        op.template reducePacket<PacketReturnType>(scratch[localid + offset], &packetAccumulator);
        scratch[localid] = op.template finalizePacket<PacketReturnType>(packetAccumulator);
      }
    }
    if (localid == 0) {
      output_ptr[itemID.get_group(0)] =
          op.finalizeBoth(op.initialize(), op.template finalizePacket<PacketReturnType>(packetAccumulator));
    }
  }

References EIGEN_UNROLL_LOOP, Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::evaluator, Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::final_output, Eigen::TensorSycl::internal::OpDefiner< Op, CoeffReturnType, Index, Vectorizable >::finalise_op(), i, Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::op, Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::rng, Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::scratch, and oomph::CumulativeTimings::start().

Referenced by Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::operator()().

compute_reduction() [2/2]

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

template<bool Vect = (Evaluator::ReducerTraits::PacketAccess & Evaluator::InputPacketAccess)>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<!Vect> Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::compute_reduction ( const cl::sycl::nd_item< 1 > &  itemID ) const

inline

                                            {
    auto output_ptr = final_output;
    Index globalid = itemID.get_global_id(0);
    Index localid = itemID.get_local_id(0);
    // vectorizable parts
    CoeffReturnType accumulator = op.initialize();
    // non vectorizable parts
    for (Index i = globalid; i < rng; i += itemID.get_global_range(0)) {
      op.reduce(evaluator.impl().coeff(i), &accumulator);
    }
    scratch[localid] = accumulator = OpDef::finalise_op(op.finalize(accumulator), rng);
 
    // reduction parts. the local size is always power of 2
    EIGEN_UNROLL_LOOP
    for (Index offset = local_range / 2; offset > 0; offset /= 2) {
      itemID.barrier(cl::sycl::access::fence_space::local_space);
      if (localid < offset) {
        op.reduce(scratch[localid + offset], &accumulator);
        scratch[localid] = op.finalize(accumulator);
      }
    }
    if (localid == 0) {
      output_ptr[itemID.get_group(0)] = op.finalize(accumulator);
    }
  }

References EIGEN_UNROLL_LOOP, Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::evaluator, Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::final_output, Eigen::TensorSycl::internal::OpDefiner< Op, CoeffReturnType, Index, Vectorizable >::finalise_op(), i, Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::op, Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::rng, and Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::scratch.

operator()()

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

void Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::operator() ( cl::sycl::nd_item< 1 >  itemID ) const

inline

{ compute_reduction(itemID); }

References Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::compute_reduction().

Member Data Documentation

evaluator

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

Evaluator Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::evaluator

Referenced by Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::compute_reduction().

final_output

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

EvaluatorPointerType Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::final_output

Referenced by Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::compute_reduction().

op

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

Op Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::op

Referenced by Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::compute_reduction().

rng

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

Index Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::rng

Referenced by Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::compute_reduction().

scratch

template<typename Evaluator , typename OpType , typename Evaluator::Index local_range>

LocalAccessor Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::scratch

Referenced by Eigen::TensorSycl::internal::FullReductionKernelFunctor< Evaluator, OpType, local_range >::compute_reduction().

---

The documentation for this class was generated from the following file:

• TensorReductionSycl.h