Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range > Struct Template Reference

#include <TensorReductionSycl.h>

Public Types
typedef cl::sycl::accessor< CoeffReturnType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local >	LocalAccessor
typedef OpDefiner< OpType, CoeffReturnType, Index, true >	OpDef
typedef OpDef::type	Op

Public Member Functions
	SecondStepFullReducer (LocalAccessor scratch_, InputAccessor aI_, OutputAccessor outAcc_, OpType op_)
void	operator() (cl::sycl::nd_item< 1 > itemID) const

Public Attributes
LocalAccessor	scratch
InputAccessor	aI
OutputAccessor	outAcc
Op	op

Member Typedef Documentation

LocalAccessor

template<typename CoeffReturnType , typename OpType , typename InputAccessor , typename OutputAccessor , typename Index , Index local_range>

typedef cl::sycl::accessor<CoeffReturnType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::LocalAccessor

Op

template<typename CoeffReturnType , typename OpType , typename InputAccessor , typename OutputAccessor , typename Index , Index local_range>

typedef OpDef::type Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::Op

OpDef

template<typename CoeffReturnType , typename OpType , typename InputAccessor , typename OutputAccessor , typename Index , Index local_range>

typedef OpDefiner<OpType, CoeffReturnType, Index, true> Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::OpDef

Constructor & Destructor Documentation

SecondStepFullReducer()

template<typename CoeffReturnType , typename OpType , typename InputAccessor , typename OutputAccessor , typename Index , Index local_range>

Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::SecondStepFullReducer ( LocalAccessor  scratch_, InputAccessor  aI_, OutputAccessor  outAcc_, OpType  op_  )

inline

      : scratch(scratch_), aI(aI_), outAcc(outAcc_), op(OpDef::get_op(op_)) {}

Member Function Documentation

operator()()

template<typename CoeffReturnType , typename OpType , typename InputAccessor , typename OutputAccessor , typename Index , Index local_range>

void Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::operator() ( cl::sycl::nd_item< 1 >  itemID ) const

inline

                                                 {
    // Our empirical research shows that the best performance will be achieved
    // when there is only one element per thread to reduce in the second step.
    // in this step the second step reduction time is almost negligible.
    // Hence, in the second step of reduction the input size is fixed to the
    // local size, thus, there is only one element read per thread. The
    // algorithm must be changed if the number of reduce per thread in the
    // second step is greater than 1. Otherwise, the result will be wrong.
    const Index localid = itemID.get_local_id(0);
    auto aInPtr = aI + localid;
    auto aOutPtr = outAcc;
    CoeffReturnType *scratchptr = scratch.get_pointer();
    CoeffReturnType accumulator = *aInPtr;
 
    scratchptr[localid] = op.finalize(accumulator);
    for (Index offset = itemID.get_local_range(0) / 2; offset > 0; offset /= 2) {
      itemID.barrier(cl::sycl::access::fence_space::local_space);
      if (localid < offset) {
        op.reduce(scratchptr[localid + offset], &accumulator);
        scratchptr[localid] = op.finalize(accumulator);
      }
    }
    if (localid == 0) *aOutPtr = op.finalize(accumulator);
  }

References Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::aI, Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::op, Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::outAcc, and Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::scratch.

Member Data Documentation

aI

template<typename CoeffReturnType , typename OpType , typename InputAccessor , typename OutputAccessor , typename Index , Index local_range>

InputAccessor Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::aI

Referenced by Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::operator()().

op

template<typename CoeffReturnType , typename OpType , typename InputAccessor , typename OutputAccessor , typename Index , Index local_range>

Op Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::op

Referenced by Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::operator()().

outAcc

template<typename CoeffReturnType , typename OpType , typename InputAccessor , typename OutputAccessor , typename Index , Index local_range>

OutputAccessor Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::outAcc

Referenced by Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::operator()().

scratch

template<typename CoeffReturnType , typename OpType , typename InputAccessor , typename OutputAccessor , typename Index , Index local_range>

LocalAccessor Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::scratch

Referenced by Eigen::TensorSycl::internal::SecondStepFullReducer< CoeffReturnType, OpType, InputAccessor, OutputAccessor, Index, local_range >::operator()().

---

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

• TensorReductionSycl.h