Eigen::TensorSycl::internal::PartialReducerLauncher< Self, Op, rt > Struct Template Reference

#include <TensorReductionSycl.h>

Public Types
typedef Self::EvaluatorPointerType	EvaluatorPointerType
typedef Self::CoeffReturnType	CoeffReturnType
typedef Self::Storage	Storage
typedef Self::Index	Index
typedef ReductionPannel< typename Self::Index, EIGEN_SYCL_LOCAL_THREAD_DIM0, EIGEN_SYCL_LOCAL_THREAD_DIM1, true >	PannelParameters
typedef PartialReductionKernel< Self, Op, PannelParameters, rt >	SyclReducerKerneType

Static Public Member Functions
static bool	run (const Self &self, const Op &reducer, const Eigen::SyclDevice &dev, EvaluatorPointerType output, Index num_coeffs_to_reduce, Index num_coeffs_to_preserve)

Member Typedef Documentation

CoeffReturnType

template<typename Self , typename Op , TensorSycl::internal::reduction_dim rt>

typedef Self::CoeffReturnType Eigen::TensorSycl::internal::PartialReducerLauncher< Self, Op, rt >::CoeffReturnType

EvaluatorPointerType

template<typename Self , typename Op , TensorSycl::internal::reduction_dim rt>

typedef Self::EvaluatorPointerType Eigen::TensorSycl::internal::PartialReducerLauncher< Self, Op, rt >::EvaluatorPointerType

Index

template<typename Self , typename Op , TensorSycl::internal::reduction_dim rt>

typedef Self::Index Eigen::TensorSycl::internal::PartialReducerLauncher< Self, Op, rt >::Index

PannelParameters

template<typename Self , typename Op , TensorSycl::internal::reduction_dim rt>

typedef ReductionPannel<typename Self::Index, EIGEN_SYCL_LOCAL_THREAD_DIM0, EIGEN_SYCL_LOCAL_THREAD_DIM1, true> Eigen::TensorSycl::internal::PartialReducerLauncher< Self, Op, rt >::PannelParameters

Storage

template<typename Self , typename Op , TensorSycl::internal::reduction_dim rt>

typedef Self::Storage Eigen::TensorSycl::internal::PartialReducerLauncher< Self, Op, rt >::Storage

SyclReducerKerneType

template<typename Self , typename Op , TensorSycl::internal::reduction_dim rt>

typedef PartialReductionKernel<Self, Op, PannelParameters, rt> Eigen::TensorSycl::internal::PartialReducerLauncher< Self, Op, rt >::SyclReducerKerneType

Member Function Documentation

run()

template<typename Self , typename Op , TensorSycl::internal::reduction_dim rt>

static bool Eigen::TensorSycl::internal::PartialReducerLauncher< Self, Op, rt >::run ( const Self &  self, const Op &  reducer, const Eigen::SyclDevice &  dev, EvaluatorPointerType  output, Index  num_coeffs_to_reduce, Index  num_coeffs_to_preserve  )

inlinestatic

                                                                            {
    Index roundUpP = roundUp(num_coeffs_to_preserve, PannelParameters::LocalThreadSizeP);
 
    // getPowerOfTwo makes sure local range is power of 2 and <=
    // maxSyclThreadPerBlock this will help us to avoid extra check on the
    // kernel
    static_assert(!((PannelParameters::LocalThreadSizeP * PannelParameters::LocalThreadSizeR) &
                    (PannelParameters::LocalThreadSizeP * PannelParameters::LocalThreadSizeR - 1)),
                  "The Local thread size must be a power of 2 for the reduction "
                  "operation");
 
    EIGEN_CONSTEXPR Index localRange = PannelParameters::LocalThreadSizeP * PannelParameters::LocalThreadSizeR;
    // In this step, we force the code not to be more than 2-step reduction:
    // Our empirical research shows that if each thread reduces at least 64
    // elements individually, we get better performance. However, this can change
    // on different platforms. In this step we force the code not to be
    // morthan step reduction: Our empirical research shows that for inner_most
    // dim reducer, it is better to have 8 group in a reduce dimension for sizes
    // > 1024 to achieve the best performance.
    const Index reductionPerThread = 64;
    Index cu = dev.getPowerOfTwo(dev.getNumSyclMultiProcessors(), true);
    const Index pNumGroups = roundUpP / PannelParameters::LocalThreadSizeP;
    Index rGroups = (cu + pNumGroups - 1) / pNumGroups;
    const Index rNumGroups = num_coeffs_to_reduce > reductionPerThread * localRange ? std::min(rGroups, localRange) : 1;
    const Index globalRange = pNumGroups * rNumGroups * localRange;
 
    EIGEN_CONSTEXPR Index scratchSize =
        PannelParameters::LocalThreadSizeR * (PannelParameters::LocalThreadSizeP + PannelParameters::BC);
    auto thread_range = cl::sycl::nd_range<1>(cl::sycl::range<1>(globalRange), cl::sycl::range<1>(localRange));
    if (rNumGroups > 1) {
      CoeffReturnType *temp_pointer = static_cast<CoeffReturnType *>(
          dev.allocate_temp(num_coeffs_to_preserve * rNumGroups * sizeof(CoeffReturnType)));
      EvaluatorPointerType temp_accessor = dev.get(temp_pointer);
      dev.template unary_kernel_launcher<CoeffReturnType, SyclReducerKerneType>(
             self, temp_accessor, thread_range, scratchSize, reducer, pNumGroups, rNumGroups, num_coeffs_to_preserve,
             num_coeffs_to_reduce)
          .wait();
      typedef SecondStepPartialReduction<CoeffReturnType, Index, EvaluatorPointerType, EvaluatorPointerType, Op>
          SecondStepPartialReductionKernel;
      dev.template unary_kernel_launcher<CoeffReturnType, SecondStepPartialReductionKernel>(
             temp_accessor, output,
             cl::sycl::nd_range<1>(cl::sycl::range<1>(pNumGroups * localRange), cl::sycl::range<1>(localRange)),
             Index(1), reducer, num_coeffs_to_preserve, rNumGroups)
          .wait();
      self.device().deallocate_temp(temp_pointer);
    } else {
      dev.template unary_kernel_launcher<CoeffReturnType, SyclReducerKerneType>(
             self, output, thread_range, scratchSize, reducer, pNumGroups, rNumGroups, num_coeffs_to_preserve,
             num_coeffs_to_reduce)
          .wait();
    }
    return false;
  }

References Eigen::TensorSycl::internal::ReductionPannel< Index, LTP, LTR, BC_ >::BC, EIGEN_CONSTEXPR, Eigen::TensorSycl::internal::ReductionPannel< Index, LTP, LTR, BC_ >::LocalThreadSizeP, Eigen::TensorSycl::internal::ReductionPannel< Index, LTP, LTR, BC_ >::LocalThreadSizeR, min, and output().

---

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

• TensorReductionSycl.h