Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt > Struct Template Reference

#include <TensorReductionSycl.h>

Public Types
typedef Evaluator::CoeffReturnType	CoeffReturnType
typedef Evaluator::EvaluatorPointerType	EvaluatorPointerType
typedef Evaluator::Index	Index
typedef OpDefiner< OpType, CoeffReturnType, Index, false >	OpDef
typedef OpDef::type	Op
typedef cl::sycl::accessor< CoeffReturnType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local >	ScratchAcc

Public Member Functions
	PartialReductionKernel (ScratchAcc scratch_, Evaluator evaluator_, EvaluatorPointerType output_accessor_, OpType op_, const Index preserve_elements_num_groups_, const Index reduce_elements_num_groups_, const Index num_coeffs_to_preserve_, const Index num_coeffs_to_reduce_)
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void	element_wise_reduce (Index globalRId, Index globalPId, CoeffReturnType &accumulator) const
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void	operator() (cl::sycl::nd_item< 1 > itemID) const

Public Attributes
ScratchAcc	scratch
Evaluator	evaluator
EvaluatorPointerType	output_accessor
Op	op
const Index	preserve_elements_num_groups
const Index	reduce_elements_num_groups
const Index	num_coeffs_to_preserve
const Index	num_coeffs_to_reduce

Member Typedef Documentation

CoeffReturnType

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

typedef Evaluator::CoeffReturnType Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::CoeffReturnType

EvaluatorPointerType

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

typedef Evaluator::EvaluatorPointerType Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::EvaluatorPointerType

Index

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

typedef Evaluator::Index Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::Index

Op

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

typedef OpDef::type Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::Op

OpDef

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

typedef OpDefiner<OpType, CoeffReturnType, Index, false> Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::OpDef

ScratchAcc

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

typedef cl::sycl::accessor<CoeffReturnType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::ScratchAcc

Constructor & Destructor Documentation

PartialReductionKernel()

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::PartialReductionKernel ( ScratchAcc  scratch_, Evaluator  evaluator_, EvaluatorPointerType  output_accessor_, OpType  op_, const Index  preserve_elements_num_groups_, const Index  reduce_elements_num_groups_, const Index  num_coeffs_to_preserve_, const Index  num_coeffs_to_reduce_  )

inline

      : scratch(scratch_),
        evaluator(evaluator_),
        output_accessor(output_accessor_),
        op(OpDef::get_op(op_)),
        preserve_elements_num_groups(preserve_elements_num_groups_),
        reduce_elements_num_groups(reduce_elements_num_groups_),
        num_coeffs_to_preserve(num_coeffs_to_preserve_),
        num_coeffs_to_reduce(num_coeffs_to_reduce_) {}

Member Function Documentation

element_wise_reduce()

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::element_wise_reduce ( Index  globalRId, Index  globalPId, CoeffReturnType &  accumulator  ) const

inline

                                                                                                     {
    if (globalPId >= num_coeffs_to_preserve) {
      return;
    }
    Index global_offset = rt == reduction_dim::outer_most ? globalPId + (globalRId * num_coeffs_to_preserve)
                                                          : globalRId + (globalPId * num_coeffs_to_reduce);
    Index localOffset = globalRId;
 
    const Index per_thread_local_stride = PannelParameters::LocalThreadSizeR * reduce_elements_num_groups;
    const Index per_thread_global_stride =
        rt == reduction_dim::outer_most ? num_coeffs_to_preserve * per_thread_local_stride : per_thread_local_stride;
    for (Index i = globalRId; i < num_coeffs_to_reduce; i += per_thread_local_stride) {
      op.reduce(evaluator.impl().coeff(global_offset), &accumulator);
      localOffset += per_thread_local_stride;
      global_offset += per_thread_global_stride;
    }
  }

References Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::evaluator, i, Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::num_coeffs_to_preserve, Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::num_coeffs_to_reduce, Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::op, Eigen::TensorSycl::internal::outer_most, and Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::reduce_elements_num_groups.

Referenced by Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::operator()().

operator()()

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::operator() ( cl::sycl::nd_item< 1 >  itemID ) const

inline

                                                                                       {
    const Index linearLocalThreadId = itemID.get_local_id(0);
    Index pLocalThreadId = rt == reduction_dim::outer_most ? linearLocalThreadId % PannelParameters::LocalThreadSizeP
                                                           : linearLocalThreadId / PannelParameters::LocalThreadSizeR;
    Index rLocalThreadId = rt == reduction_dim::outer_most ? linearLocalThreadId / PannelParameters::LocalThreadSizeP
                                                           : linearLocalThreadId % PannelParameters::LocalThreadSizeR;
    const Index pGroupId = rt == reduction_dim::outer_most ? itemID.get_group(0) % preserve_elements_num_groups
                                                           : itemID.get_group(0) / reduce_elements_num_groups;
    const Index rGroupId = rt == reduction_dim::outer_most ? itemID.get_group(0) / preserve_elements_num_groups
                                                           : itemID.get_group(0) % reduce_elements_num_groups;
 
    Index globalPId = pGroupId * PannelParameters::LocalThreadSizeP + pLocalThreadId;
    const Index globalRId = rGroupId * PannelParameters::LocalThreadSizeR + rLocalThreadId;
    CoeffReturnType *scratchPtr = scratch.get_pointer();
    auto outPtr = output_accessor + (reduce_elements_num_groups > 1 ? rGroupId * num_coeffs_to_preserve : 0);
    CoeffReturnType accumulator = op.initialize();
 
    element_wise_reduce(globalRId, globalPId, accumulator);
 
    accumulator = OpDef::finalise_op(op.finalize(accumulator), num_coeffs_to_reduce);
    scratchPtr[pLocalThreadId + rLocalThreadId * (PannelParameters::LocalThreadSizeP + PannelParameters::BC)] =
        accumulator;
    if (rt == reduction_dim::inner_most) {
      pLocalThreadId = linearLocalThreadId % PannelParameters::LocalThreadSizeP;
      rLocalThreadId = linearLocalThreadId / PannelParameters::LocalThreadSizeP;
      globalPId = pGroupId * PannelParameters::LocalThreadSizeP + pLocalThreadId;
    }
 
    /* Apply the reduction operation between the current local
     * id and the one on the other half of the vector. */
    auto out_scratch_ptr =
        scratchPtr + (pLocalThreadId + (rLocalThreadId * (PannelParameters::LocalThreadSizeP + PannelParameters::BC)));
    itemID.barrier(cl::sycl::access::fence_space::local_space);
    if (rt == reduction_dim::inner_most) {
      accumulator = *out_scratch_ptr;
    }
    // The Local LocalThreadSizeR is always power of 2
    EIGEN_UNROLL_LOOP
    for (Index offset = PannelParameters::LocalThreadSizeR >> 1; offset > 0; offset >>= 1) {
      if (rLocalThreadId < offset) {
        op.reduce(out_scratch_ptr[(PannelParameters::LocalThreadSizeP + PannelParameters::BC) * offset], &accumulator);
        // The result has already been divided for mean reducer in the
        // previous reduction so no need to divide furthermore
        *out_scratch_ptr = op.finalize(accumulator);
      }
      /* All threads collectively read from global memory into local.
       * The barrier ensures all threads' IO is resolved before
       * execution continues (strictly speaking, all threads within
       * a single work-group - there is no co-ordination between
       * work-groups, only work-items). */
      itemID.barrier(cl::sycl::access::fence_space::local_space);
    }
 
    if (rLocalThreadId == 0 && (globalPId < num_coeffs_to_preserve)) {
      outPtr[globalPId] = op.finalize(accumulator);
    }
  }

References EIGEN_UNROLL_LOOP, Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::element_wise_reduce(), Eigen::TensorSycl::internal::OpDefiner< Op, CoeffReturnType, Index, Vectorizable >::finalise_op(), Eigen::TensorSycl::internal::inner_most, Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::num_coeffs_to_preserve, Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::num_coeffs_to_reduce, Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::op, Eigen::TensorSycl::internal::outer_most, Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::output_accessor, Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::preserve_elements_num_groups, Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::reduce_elements_num_groups, and Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::scratch.

Member Data Documentation

evaluator

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

Evaluator Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::evaluator

Referenced by Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::element_wise_reduce().

num_coeffs_to_preserve

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

const Index Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::num_coeffs_to_preserve

Referenced by Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::element_wise_reduce(), and Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::operator()().

num_coeffs_to_reduce

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

const Index Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::num_coeffs_to_reduce

Referenced by Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::element_wise_reduce(), and Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::operator()().

op

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

Op Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::op

Referenced by Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::element_wise_reduce(), and Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::operator()().

output_accessor

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

EvaluatorPointerType Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::output_accessor

Referenced by Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::operator()().

preserve_elements_num_groups

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

const Index Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::preserve_elements_num_groups

Referenced by Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::operator()().

reduce_elements_num_groups

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

const Index Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::reduce_elements_num_groups

Referenced by Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::element_wise_reduce(), and Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::operator()().

scratch

template<typename Evaluator , typename OpType , typename PannelParameters , reduction_dim rt>

ScratchAcc Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::scratch

Referenced by Eigen::TensorSycl::internal::PartialReductionKernel< Evaluator, OpType, PannelParameters, rt >::operator()().

---

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

• TensorReductionSycl.h