Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index > 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	ScanAdjustmentKernelFunctor (LocalAccessor, InAccessor in_accessor_, OutAccessor out_accessor_, const ScanParameters< Index > scanParameters_, Op accumulator_)
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void	operator() (cl::sycl::nd_item< 1 > itemID) const

Public Attributes
InAccessor	in_ptr
OutAccessor	out_ptr
const ScanParameters< Index >	scanParameters
Op	accumulator

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

Member Typedef Documentation

LocalAccessor

template<typename CoeffReturnType , typename InAccessor , typename OutAccessor , typename Op , typename Index >

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

Constructor & Destructor Documentation

ScanAdjustmentKernelFunctor()

template<typename CoeffReturnType , typename InAccessor , typename OutAccessor , typename Op , typename Index >

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::ScanAdjustmentKernelFunctor ( LocalAccessor  , InAccessor  in_accessor_, OutAccessor  out_accessor_, const ScanParameters< Index >  scanParameters_, Op  accumulator_  )

inline

      : in_ptr(in_accessor_), out_ptr(out_accessor_), scanParameters(scanParameters_), accumulator(accumulator_) {}

Member Function Documentation

operator()()

template<typename CoeffReturnType , typename InAccessor , typename OutAccessor , typename Op , typename Index >

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::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;
 
      // 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;
      const Index block_size = scanParameters.elements_per_group / scanParameters.elements_per_block;
      const Index in_id = (panel_id * block_size * scanParameters.non_scan_size) + (group_id * block_size) + block_id;
      CoeffReturnType adjust_val = in_ptr[in_id];
 
      Index next_elements = 0;
      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)) {
          CoeffReturnType accum = adjust_val;
          accumulator.reduce(out_ptr[global_id], &accum);
          out_ptr[global_id] = accumulator.finalize(accum);
        }
        next_elements += scanParameters.scan_stride;
      }
    }
  }

References Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::accumulator, Eigen::TensorSycl::internal::ScanParameters< index_t >::block_threads, EIGEN_UNROLL_LOOP, Eigen::TensorSycl::internal::ScanParameters< index_t >::elements_per_block, Eigen::TensorSycl::internal::ScanParameters< index_t >::elements_per_group, Eigen::TensorSycl::internal::ScanParameters< index_t >::group_threads, i, Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::in_ptr, 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::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::out_ptr, Eigen::TensorSycl::internal::ScanParameters< index_t >::panel_threads, Eigen::TensorSycl::internal::ScanParameters< index_t >::scan_size, Eigen::TensorSycl::internal::ScanParameters< index_t >::scan_stride, Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::scanParameters, tmp, and Eigen::TensorSycl::internal::ScanParameters< index_t >::total_size.

Member Data Documentation

accumulator

template<typename CoeffReturnType , typename InAccessor , typename OutAccessor , typename Op , typename Index >

Op Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::accumulator

Referenced by Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::operator()().

in_ptr

template<typename CoeffReturnType , typename InAccessor , typename OutAccessor , typename Op , typename Index >

InAccessor Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::in_ptr

Referenced by Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::operator()().

out_ptr

template<typename CoeffReturnType , typename InAccessor , typename OutAccessor , typename Op , typename Index >

OutAccessor Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::out_ptr

Referenced by Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::operator()().

PacketSize

template<typename CoeffReturnType , typename InAccessor , typename OutAccessor , typename Op , typename Index >

EIGEN_CONSTEXPR int Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::PacketSize = ScanParameters<Index>::ScanPerThread / 2

static

scanParameters

template<typename CoeffReturnType , typename InAccessor , typename OutAccessor , typename Op , typename Index >

const ScanParameters<Index> Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::scanParameters

Referenced by Eigen::TensorSycl::internal::ScanAdjustmentKernelFunctor< CoeffReturnType, InAccessor, OutAccessor, Op, Index >::operator()().

---

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

• TensorScanSycl.h