Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp > Struct Template Reference

#include <TensorScanSycl.h>

Public Types
typedef cl::sycl::accessor< CoeffReturnType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local >	LocalAccessor

Public Member Functions
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE	ScanKernelFunctor (LocalAccessor scratch_, const Evaluator dev_eval_, OutAccessor out_accessor_, OutAccessor temp_accessor_, const ScanParameters< Index > scanParameters_, Op accumulator_, const bool inclusive_)
template<scan_step sst = stp, typename Input >
std::enable_if_t< sst==scan_step::first, CoeffReturnType > EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE	read (const Input &inpt, Index global_id) const
template<scan_step sst = stp, typename Input >
std::enable_if_t< sst !=scan_step::first, CoeffReturnType > EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE	read (const Input &inpt, Index global_id) const
template<scan_step sst = stp, typename InclusiveOp >
std::enable_if_t< sst==scan_step::first > EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE	first_step_inclusive_Operation (InclusiveOp inclusive_op) const
template<scan_step sst = stp, typename InclusiveOp >
std::enable_if_t< sst !=scan_step::first > EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE	first_step_inclusive_Operation (InclusiveOp) const
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void	operator() (cl::sycl::nd_item< 1 > itemID) const

Public Attributes
LocalAccessor	scratch
Evaluator	dev_eval
OutAccessor	out_ptr
OutAccessor	tmp_ptr
const ScanParameters< Index >	scanParameters
Op	accumulator
const bool	inclusive

Static Public Attributes
static EIGEN_CONSTEXPR int	PacketSize = ScanParameters<Index>::ScanPerThread / 2

Member Typedef Documentation

LocalAccessor

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

typedef cl::sycl::accessor<CoeffReturnType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::LocalAccessor

Constructor & Destructor Documentation

ScanKernelFunctor()

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::ScanKernelFunctor ( LocalAccessor  scratch_, const Evaluator  dev_eval_, OutAccessor  out_accessor_, OutAccessor  temp_accessor_, const ScanParameters< Index >  scanParameters_, Op  accumulator_, const bool  inclusive_  )

inline

      : scratch(scratch_),
        dev_eval(dev_eval_),
        out_ptr(out_accessor_),
        tmp_ptr(temp_accessor_),
        scanParameters(scanParameters_),
        accumulator(accumulator_),
        inclusive(inclusive_) {}

Member Function Documentation

first_step_inclusive_Operation() [1/2]

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

template<scan_step sst = stp, typename InclusiveOp >

std::enable_if_t<sst == scan_step::first> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::first_step_inclusive_Operation ( InclusiveOp  inclusive_op ) const

inline

                                      {
    inclusive_op();
  }

Referenced by Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::operator()().

first_step_inclusive_Operation() [2/2]

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

template<scan_step sst = stp, typename InclusiveOp >

std::enable_if_t<sst != scan_step::first> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::first_step_inclusive_Operation ( InclusiveOp  ) const

inline

                         {}

operator()()

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::operator() ( cl::sycl::nd_item< 1 >  itemID ) const

inline

                                                                                       {
    for (Index loop_offset = 0; loop_offset < scanParameters.loop_range; loop_offset++) {
      Index data_offset = (itemID.get_global_id(0) + (itemID.get_global_range(0) * loop_offset));
      Index tmp = data_offset % scanParameters.panel_threads;
      const Index panel_id = data_offset / scanParameters.panel_threads;
      const Index group_id = tmp / scanParameters.group_threads;
      tmp = tmp % scanParameters.group_threads;
      const Index block_id = tmp / scanParameters.block_threads;
      const Index local_id = tmp % scanParameters.block_threads;
      // we put one element per packet in scratch_mem
      const Index scratch_stride = scanParameters.elements_per_block / PacketSize;
      const Index scratch_offset = (itemID.get_local_id(0) / scanParameters.block_threads) * scratch_stride;
      CoeffReturnType private_scan[ScanParameters<Index>::ScanPerThread];
      CoeffReturnType inclusive_scan;
      // the actual panel size is scan_size * non_scan_size.
      // elements_per_panel is roundup to power of 2 for binary tree
      const Index panel_offset = panel_id * scanParameters.scan_size * scanParameters.non_scan_size;
      const Index group_offset = group_id * scanParameters.non_scan_stride;
      // This will be effective when the size is bigger than elements_per_block
      const Index block_offset = block_id * scanParameters.elements_per_block * scanParameters.scan_stride;
      const Index thread_offset = (ScanParameters<Index>::ScanPerThread * local_id * scanParameters.scan_stride);
      const Index global_offset = panel_offset + group_offset + block_offset + thread_offset;
      Index next_elements = 0;
      EIGEN_UNROLL_LOOP
      for (int i = 0; i < ScanParameters<Index>::ScanPerThread; i++) {
        Index global_id = global_offset + next_elements;
        private_scan[i] = ((((block_id * scanParameters.elements_per_block) +
                             (ScanParameters<Index>::ScanPerThread * local_id) + i) < scanParameters.scan_size) &&
                           (global_id < scanParameters.total_size))
                              ? read(dev_eval, global_id)
                              : accumulator.initialize();
        next_elements += scanParameters.scan_stride;
      }
      first_step_inclusive_Operation([&]() EIGEN_DEVICE_FUNC {
        if (inclusive) {
          inclusive_scan = private_scan[ScanParameters<Index>::ScanPerThread - 1];
        }
      });
      // This for loop must be 2
      EIGEN_UNROLL_LOOP
      for (int packetIndex = 0; packetIndex < ScanParameters<Index>::ScanPerThread; packetIndex += PacketSize) {
        Index private_offset = 1;
        // build sum in place up the tree
        EIGEN_UNROLL_LOOP
        for (Index d = PacketSize >> 1; d > 0; d >>= 1) {
          EIGEN_UNROLL_LOOP
          for (Index l = 0; l < d; l++) {
            Index ai = private_offset * (2 * l + 1) - 1 + packetIndex;
            Index bi = private_offset * (2 * l + 2) - 1 + packetIndex;
            CoeffReturnType accum = accumulator.initialize();
            accumulator.reduce(private_scan[ai], &accum);
            accumulator.reduce(private_scan[bi], &accum);
            private_scan[bi] = accumulator.finalize(accum);
          }
          private_offset *= 2;
        }
        scratch[2 * local_id + (packetIndex / PacketSize) + scratch_offset] =
            private_scan[PacketSize - 1 + packetIndex];
        private_scan[PacketSize - 1 + packetIndex] = accumulator.initialize();
        // traverse down tree & build scan
        EIGEN_UNROLL_LOOP
        for (Index d = 1; d < PacketSize; d *= 2) {
          private_offset >>= 1;
          EIGEN_UNROLL_LOOP
          for (Index l = 0; l < d; l++) {
            Index ai = private_offset * (2 * l + 1) - 1 + packetIndex;
            Index bi = private_offset * (2 * l + 2) - 1 + packetIndex;
            CoeffReturnType accum = accumulator.initialize();
            accumulator.reduce(private_scan[ai], &accum);
            accumulator.reduce(private_scan[bi], &accum);
            private_scan[ai] = private_scan[bi];
            private_scan[bi] = accumulator.finalize(accum);
          }
        }
      }
 
      Index offset = 1;
      // build sum in place up the tree
      for (Index d = scratch_stride >> 1; d > 0; d >>= 1) {
        // Synchronise
        itemID.barrier(cl::sycl::access::fence_space::local_space);
        if (local_id < d) {
          Index ai = offset * (2 * local_id + 1) - 1 + scratch_offset;
          Index bi = offset * (2 * local_id + 2) - 1 + scratch_offset;
          CoeffReturnType accum = accumulator.initialize();
          accumulator.reduce(scratch[ai], &accum);
          accumulator.reduce(scratch[bi], &accum);
          scratch[bi] = accumulator.finalize(accum);
        }
        offset *= 2;
      }
      // Synchronise
      itemID.barrier(cl::sycl::access::fence_space::local_space);
      // next step optimisation
      if (local_id == 0) {
        if (((scanParameters.elements_per_group / scanParameters.elements_per_block) > 1)) {
          const Index temp_id = panel_id * (scanParameters.elements_per_group / scanParameters.elements_per_block) *
                                    scanParameters.non_scan_size +
                                group_id * (scanParameters.elements_per_group / scanParameters.elements_per_block) +
                                block_id;
          tmp_ptr[temp_id] = scratch[scratch_stride - 1 + scratch_offset];
        }
        // clear the last element
        scratch[scratch_stride - 1 + scratch_offset] = accumulator.initialize();
      }
      // traverse down tree & build scan
      for (Index d = 1; d < scratch_stride; d *= 2) {
        offset >>= 1;
        // Synchronise
        itemID.barrier(cl::sycl::access::fence_space::local_space);
        if (local_id < d) {
          Index ai = offset * (2 * local_id + 1) - 1 + scratch_offset;
          Index bi = offset * (2 * local_id + 2) - 1 + scratch_offset;
          CoeffReturnType accum = accumulator.initialize();
          accumulator.reduce(scratch[ai], &accum);
          accumulator.reduce(scratch[bi], &accum);
          scratch[ai] = scratch[bi];
          scratch[bi] = accumulator.finalize(accum);
        }
      }
      // Synchronise
      itemID.barrier(cl::sycl::access::fence_space::local_space);
      // This for loop must be 2
      EIGEN_UNROLL_LOOP
      for (int packetIndex = 0; packetIndex < ScanParameters<Index>::ScanPerThread; packetIndex += PacketSize) {
        EIGEN_UNROLL_LOOP
        for (Index i = 0; i < PacketSize; i++) {
          CoeffReturnType accum = private_scan[packetIndex + i];
          accumulator.reduce(scratch[2 * local_id + (packetIndex / PacketSize) + scratch_offset], &accum);
          private_scan[packetIndex + i] = accumulator.finalize(accum);
        }
      }
      first_step_inclusive_Operation([&]() EIGEN_DEVICE_FUNC {
        if (inclusive) {
          accumulator.reduce(private_scan[ScanParameters<Index>::ScanPerThread - 1], &inclusive_scan);
          private_scan[0] = accumulator.finalize(inclusive_scan);
        }
      });
      next_elements = 0;
      // right the first set of private param
      EIGEN_UNROLL_LOOP
      for (Index i = 0; i < ScanParameters<Index>::ScanPerThread; i++) {
        Index global_id = global_offset + next_elements;
        if ((((block_id * scanParameters.elements_per_block) + (ScanParameters<Index>::ScanPerThread * local_id) + i) <
             scanParameters.scan_size) &&
            (global_id < scanParameters.total_size)) {
          Index private_id = (i * !inclusive) + (((i + 1) % ScanParameters<Index>::ScanPerThread) * (inclusive));
          out_ptr[global_id] = private_scan[private_id];
        }
        next_elements += scanParameters.scan_stride;
      }
    }  // end for loop
  }

References Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::accumulator, Eigen::TensorSycl::internal::ScanParameters< index_t >::block_threads, Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::dev_eval, EIGEN_DEVICE_FUNC, EIGEN_UNROLL_LOOP, Eigen::TensorSycl::internal::ScanParameters< index_t >::elements_per_block, Eigen::TensorSycl::internal::ScanParameters< index_t >::elements_per_group, Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::first_step_inclusive_Operation(), Eigen::TensorSycl::internal::ScanParameters< index_t >::group_threads, i, Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::inclusive, Eigen::TensorSycl::internal::ScanParameters< index_t >::loop_range, Eigen::TensorSycl::internal::ScanParameters< index_t >::non_scan_size, Eigen::TensorSycl::internal::ScanParameters< index_t >::non_scan_stride, Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::out_ptr, Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::PacketSize, Eigen::TensorSycl::internal::ScanParameters< index_t >::panel_threads, Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::read(), Eigen::TensorSycl::internal::ScanParameters< index_t >::scan_size, Eigen::TensorSycl::internal::ScanParameters< index_t >::scan_stride, Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::scanParameters, Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::scratch, tmp, Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::tmp_ptr, and Eigen::TensorSycl::internal::ScanParameters< index_t >::total_size.

read() [1/2]

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

template<scan_step sst = stp, typename Input >

std::enable_if_t<sst == scan_step::first, CoeffReturnType> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::read ( const Input &  inpt, Index  global_id  ) const

inline

                                                {
    return inpt.coeff(global_id);
  }

Referenced by Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::operator()().

read() [2/2]

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

template<scan_step sst = stp, typename Input >

std::enable_if_t<sst != scan_step::first, CoeffReturnType> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::read ( const Input &  inpt, Index  global_id  ) const

inline

                                                {
    return inpt[global_id];
  }

Member Data Documentation

accumulator

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

Op Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::accumulator

Referenced by Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::operator()().

dev_eval

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

Evaluator Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::dev_eval

Referenced by Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::operator()().

inclusive

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

const bool Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::inclusive

Referenced by Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::operator()().

out_ptr

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

OutAccessor Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::out_ptr

Referenced by Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::operator()().

PacketSize

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

EIGEN_CONSTEXPR int Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::PacketSize = ScanParameters<Index>::ScanPerThread / 2

static

Referenced by Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::operator()().

scanParameters

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

const ScanParameters<Index> Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::scanParameters

Referenced by Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::operator()().

scratch

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

LocalAccessor Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::scratch

Referenced by Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::operator()().

tmp_ptr

template<typename Evaluator , typename CoeffReturnType , typename OutAccessor , typename Op , typename Index , scan_step stp>

OutAccessor Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::tmp_ptr

Referenced by Eigen::TensorSycl::internal::ScanKernelFunctor< Evaluator, CoeffReturnType, OutAccessor, Op, Index, stp >::operator()().

---

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

• TensorScanSycl.h