| // This file is part of Eigen, a lightweight C++ template library |
| // for linear algebra. |
| // |
| // Mehdi Goli Codeplay Software Ltd. |
| // Ralph Potter Codeplay Software Ltd. |
| // Luke Iwanski Codeplay Software Ltd. |
| // Contact: <eigen@codeplay.com> |
| // |
| // This Source Code Form is subject to the terms of the Mozilla |
| // Public License v. 2.0. If a copy of the MPL was not distributed |
| // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. |
| |
| /***************************************************************** |
| * TensorReductionSycl.h |
| * |
| * \brief: |
| * This is the specialization of the reduction operation. Two phase reduction approach |
| * is used since the GPU does not have Global Synchronization for global memory among |
| * different work-group/thread block. To solve the problem, we need to create two kernels |
| * to reduce the data, where the first kernel reduce the data locally and each local |
| * workgroup/thread-block save the input data into global memory. In the second phase (global reduction) |
| * one work-group uses one work-group/thread-block to reduces the intermediate data into one single element. |
| * Here is an NVIDIA presentation explaining the optimized two phase reduction algorithm on GPU: |
| * https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf |
| * |
| *****************************************************************/ |
| |
| #ifndef UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSOR_REDUCTION_SYCL_HPP |
| #define UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSOR_REDUCTION_SYCL_HPP |
| // IWYU pragma: private |
| #include "./InternalHeaderCheck.h" |
| |
| namespace Eigen { |
| namespace TensorSycl { |
| namespace internal { |
| |
| template <typename Op, typename CoeffReturnType, typename Index, bool Vectorizable> |
| struct OpDefiner { |
| typedef typename Vectorise<CoeffReturnType, Eigen::SyclDevice, Vectorizable>::PacketReturnType PacketReturnType; |
| typedef Op type; |
| static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE type get_op(Op &op) { return op; } |
| |
| static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType finalise_op(const PacketReturnType &accumulator, |
| const Index &) { |
| return accumulator; |
| } |
| }; |
| |
| template <typename CoeffReturnType, typename Index> |
| struct OpDefiner<Eigen::internal::MeanReducer<CoeffReturnType>, CoeffReturnType, Index, false> { |
| typedef Eigen::internal::SumReducer<CoeffReturnType> type; |
| static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE type get_op(Eigen::internal::MeanReducer<CoeffReturnType> &) { |
| return type(); |
| } |
| |
| static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType finalise_op(const CoeffReturnType &accumulator, |
| const Index &scale) { |
| ::Eigen::internal::scalar_quotient_op<CoeffReturnType> quotient_op; |
| return quotient_op(accumulator, CoeffReturnType(scale)); |
| } |
| }; |
| |
| template <typename CoeffReturnType, typename Index> |
| struct OpDefiner<Eigen::internal::MeanReducer<CoeffReturnType>, CoeffReturnType, Index, true> { |
| typedef typename Vectorise<CoeffReturnType, Eigen::SyclDevice, true>::PacketReturnType PacketReturnType; |
| typedef Eigen::internal::SumReducer<CoeffReturnType> type; |
| static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE type get_op(Eigen::internal::MeanReducer<CoeffReturnType> &) { |
| return type(); |
| } |
| |
| static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType finalise_op(const PacketReturnType &accumulator, |
| const Index &scale) { |
| return ::Eigen::internal::pdiv(accumulator, ::Eigen::internal::pset1<PacketReturnType>(CoeffReturnType(scale))); |
| } |
| }; |
| |
| template <typename CoeffReturnType, typename OpType, typename InputAccessor, typename OutputAccessor, typename Index, |
| Index local_range> |
| struct SecondStepFullReducer { |
| 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 typename OpDef::type Op; |
| LocalAccessor scratch; |
| InputAccessor aI; |
| OutputAccessor outAcc; |
| Op op; |
| SecondStepFullReducer(LocalAccessor scratch_, InputAccessor aI_, OutputAccessor outAcc_, OpType op_) |
| : scratch(scratch_), aI(aI_), outAcc(outAcc_), op(OpDef::get_op(op_)) {} |
| |
| void operator()(cl::sycl::nd_item<1> itemID) const { |
| // 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); |
| } |
| }; |
| |
| // Full reduction first phase. In this version the vectorization is true and the reduction accept |
| // any generic reducerOp e.g( max, min, sum, mean, iamax, iamin, etc ). |
| template <typename Evaluator, typename OpType, typename Evaluator::Index local_range> |
| class FullReductionKernelFunctor { |
| public: |
| typedef typename Evaluator::CoeffReturnType CoeffReturnType; |
| typedef typename Evaluator::Index Index; |
| typedef OpDefiner<OpType, typename Evaluator::CoeffReturnType, Index, |
| (Evaluator::ReducerTraits::PacketAccess & Evaluator::InputPacketAccess)> |
| OpDef; |
| |
| typedef typename OpDef::type Op; |
| typedef typename Evaluator::EvaluatorPointerType EvaluatorPointerType; |
| typedef typename Evaluator::PacketReturnType PacketReturnType; |
| typedef std::conditional_t<(Evaluator::ReducerTraits::PacketAccess & Evaluator::InputPacketAccess), PacketReturnType, |
| CoeffReturnType> |
| OutType; |
| typedef cl::sycl::accessor<OutType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> |
| LocalAccessor; |
| LocalAccessor scratch; |
| Evaluator evaluator; |
| EvaluatorPointerType final_output; |
| Index rng; |
| Op op; |
| |
| FullReductionKernelFunctor(LocalAccessor scratch_, Evaluator evaluator_, EvaluatorPointerType final_output_, |
| Index rng_, OpType op_) |
| : scratch(scratch_), evaluator(evaluator_), final_output(final_output_), rng(rng_), op(OpDef::get_op(op_)) {} |
| |
| void operator()(cl::sycl::nd_item<1> itemID) const { compute_reduction(itemID); } |
| |
| template <bool Vect = (Evaluator::ReducerTraits::PacketAccess & Evaluator::InputPacketAccess)> |
| EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<Vect> compute_reduction( |
| const cl::sycl::nd_item<1> &itemID) const { |
| auto output_ptr = final_output; |
| Index VectorizedRange = (rng / Evaluator::PacketSize) * Evaluator::PacketSize; |
| Index globalid = itemID.get_global_id(0); |
| Index localid = itemID.get_local_id(0); |
| Index step = Evaluator::PacketSize * itemID.get_global_range(0); |
| Index start = Evaluator::PacketSize * globalid; |
| // vectorizable parts |
| PacketReturnType packetAccumulator = op.template initializePacket<PacketReturnType>(); |
| for (Index i = start; i < VectorizedRange; i += step) { |
| op.template reducePacket<PacketReturnType>(evaluator.impl().template packet<Unaligned>(i), &packetAccumulator); |
| } |
| globalid += VectorizedRange; |
| // non vectorizable parts |
| for (Index i = globalid; i < rng; i += itemID.get_global_range(0)) { |
| op.template reducePacket<PacketReturnType>( |
| ::Eigen::TensorSycl::internal::PacketWrapper<PacketReturnType, Evaluator::PacketSize>::convert_to_packet_type( |
| evaluator.impl().coeff(i), op.initialize()), |
| &packetAccumulator); |
| } |
| scratch[localid] = packetAccumulator = |
| OpDef::finalise_op(op.template finalizePacket<PacketReturnType>(packetAccumulator), rng); |
| // reduction parts // Local size is always power of 2 |
| EIGEN_UNROLL_LOOP |
| for (Index offset = local_range / 2; offset > 0; offset /= 2) { |
| itemID.barrier(cl::sycl::access::fence_space::local_space); |
| if (localid < offset) { |
| op.template reducePacket<PacketReturnType>(scratch[localid + offset], &packetAccumulator); |
| scratch[localid] = op.template finalizePacket<PacketReturnType>(packetAccumulator); |
| } |
| } |
| if (localid == 0) { |
| output_ptr[itemID.get_group(0)] = |
| op.finalizeBoth(op.initialize(), op.template finalizePacket<PacketReturnType>(packetAccumulator)); |
| } |
| } |
| |
| template <bool Vect = (Evaluator::ReducerTraits::PacketAccess & Evaluator::InputPacketAccess)> |
| EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<!Vect> compute_reduction( |
| const cl::sycl::nd_item<1> &itemID) const { |
| auto output_ptr = final_output; |
| Index globalid = itemID.get_global_id(0); |
| Index localid = itemID.get_local_id(0); |
| // vectorizable parts |
| CoeffReturnType accumulator = op.initialize(); |
| // non vectorizable parts |
| for (Index i = globalid; i < rng; i += itemID.get_global_range(0)) { |
| op.reduce(evaluator.impl().coeff(i), &accumulator); |
| } |
| scratch[localid] = accumulator = OpDef::finalise_op(op.finalize(accumulator), rng); |
| |
| // reduction parts. the local size is always power of 2 |
| EIGEN_UNROLL_LOOP |
| for (Index offset = local_range / 2; offset > 0; offset /= 2) { |
| itemID.barrier(cl::sycl::access::fence_space::local_space); |
| if (localid < offset) { |
| op.reduce(scratch[localid + offset], &accumulator); |
| scratch[localid] = op.finalize(accumulator); |
| } |
| } |
| if (localid == 0) { |
| output_ptr[itemID.get_group(0)] = op.finalize(accumulator); |
| } |
| } |
| }; |
| |
| template <typename Evaluator, typename OpType> |
| class GenericNondeterministicReducer { |
| public: |
| typedef typename Evaluator::CoeffReturnType CoeffReturnType; |
| typedef typename Evaluator::EvaluatorPointerType EvaluatorPointerType; |
| typedef typename Evaluator::Index Index; |
| typedef OpDefiner<OpType, CoeffReturnType, Index, false> OpDef; |
| typedef typename OpDef::type Op; |
| template <typename Scratch> |
| GenericNondeterministicReducer(Scratch, Evaluator evaluator_, EvaluatorPointerType output_accessor_, OpType functor_, |
| Index range_, Index num_values_to_reduce_) |
| : evaluator(evaluator_), |
| output_accessor(output_accessor_), |
| functor(OpDef::get_op(functor_)), |
| range(range_), |
| num_values_to_reduce(num_values_to_reduce_) {} |
| |
| void operator()(cl::sycl::nd_item<1> itemID) const { |
| // This is to bypass the statefull condition in Eigen meanReducer |
| Op non_const_functor; |
| std::memcpy(&non_const_functor, &functor, sizeof(Op)); |
| auto output_accessor_ptr = output_accessor; |
| Index globalid = static_cast<Index>(itemID.get_global_linear_id()); |
| if (globalid < range) { |
| CoeffReturnType accum = functor.initialize(); |
| Eigen::internal::GenericDimReducer<Evaluator::NumReducedDims - 1, Evaluator, Op>::reduce( |
| evaluator, evaluator.firstInput(globalid), non_const_functor, &accum); |
| output_accessor_ptr[globalid] = OpDef::finalise_op(functor.finalize(accum), num_values_to_reduce); |
| } |
| } |
| |
| private: |
| Evaluator evaluator; |
| EvaluatorPointerType output_accessor; |
| Op functor; |
| Index range; |
| Index num_values_to_reduce; |
| }; |
| |
| enum class reduction_dim { inner_most, outer_most }; |
| // default is preserver |
| template <typename Evaluator, typename OpType, typename PannelParameters, reduction_dim rt> |
| struct PartialReductionKernel { |
| typedef typename Evaluator::CoeffReturnType CoeffReturnType; |
| typedef typename Evaluator::EvaluatorPointerType EvaluatorPointerType; |
| typedef typename Evaluator::Index Index; |
| typedef OpDefiner<OpType, CoeffReturnType, Index, false> OpDef; |
| typedef typename OpDef::type Op; |
| typedef cl::sycl::accessor<CoeffReturnType, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> |
| ScratchAcc; |
| 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; |
| |
| 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_) |
| : 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_) {} |
| |
| EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void element_wise_reduce(Index globalRId, Index globalPId, |
| CoeffReturnType &accumulator) const { |
| 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; |
| } |
| } |
| EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void operator()(cl::sycl::nd_item<1> itemID) const { |
| 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); |
| } |
| } |
| }; |
| |
| template <typename OutScalar, typename Index, typename InputAccessor, typename OutputAccessor, typename OpType> |
| struct SecondStepPartialReduction { |
| typedef OpDefiner<OpType, OutScalar, Index, false> OpDef; |
| typedef typename OpDef::type Op; |
| typedef cl::sycl::accessor<OutScalar, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> |
| ScratchAccessor; |
| InputAccessor input_accessor; |
| OutputAccessor output_accessor; |
| Op op; |
| const Index num_coeffs_to_preserve; |
| const Index num_coeffs_to_reduce; |
| |
| EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE SecondStepPartialReduction(ScratchAccessor, InputAccessor input_accessor_, |
| OutputAccessor output_accessor_, OpType op_, |
| const Index num_coeffs_to_preserve_, |
| const Index num_coeffs_to_reduce_) |
| : input_accessor(input_accessor_), |
| output_accessor(output_accessor_), |
| op(OpDef::get_op(op_)), |
| num_coeffs_to_preserve(num_coeffs_to_preserve_), |
| num_coeffs_to_reduce(num_coeffs_to_reduce_) {} |
| |
| EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void operator()(cl::sycl::nd_item<1> itemID) const { |
| const Index globalId = itemID.get_global_id(0); |
| |
| if (globalId >= num_coeffs_to_preserve) return; |
| |
| auto in_ptr = input_accessor + globalId; |
| |
| OutScalar accumulator = op.initialize(); |
| // num_coeffs_to_reduce is not bigger that 256 |
| for (Index i = 0; i < num_coeffs_to_reduce; i++) { |
| op.reduce(*in_ptr, &accumulator); |
| in_ptr += num_coeffs_to_preserve; |
| } |
| output_accessor[globalId] = op.finalize(accumulator); |
| } |
| }; // namespace internal |
| |
| template <typename Index, Index LTP, Index LTR, bool BC_> |
| struct ReductionPannel { |
| static EIGEN_CONSTEXPR Index LocalThreadSizeP = LTP; |
| static EIGEN_CONSTEXPR Index LocalThreadSizeR = LTR; |
| static EIGEN_CONSTEXPR bool BC = BC_; |
| }; |
| |
| template <typename Self, typename Op, TensorSycl::internal::reduction_dim rt> |
| struct PartialReducerLauncher { |
| typedef typename Self::EvaluatorPointerType EvaluatorPointerType; |
| typedef typename Self::CoeffReturnType CoeffReturnType; |
| typedef typename Self::Storage Storage; |
| typedef typename 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 bool run(const Self &self, const Op &reducer, const Eigen::SyclDevice &dev, EvaluatorPointerType output, |
| Index num_coeffs_to_reduce, Index num_coeffs_to_preserve) { |
| 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 |
| // elemnts 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; |
| } |
| }; |
| } // namespace internal |
| } // namespace TensorSycl |
| |
| namespace internal { |
| |
| template <typename Self, typename Op, bool Vectorizable> |
| struct FullReducer<Self, Op, Eigen::SyclDevice, Vectorizable> { |
| typedef typename Self::CoeffReturnType CoeffReturnType; |
| typedef typename Self::EvaluatorPointerType EvaluatorPointerType; |
| static EIGEN_CONSTEXPR bool HasOptimizedImplementation = true; |
| static EIGEN_CONSTEXPR int PacketSize = Self::PacketAccess ? Self::PacketSize : 1; |
| static void run(const Self &self, Op &reducer, const Eigen::SyclDevice &dev, EvaluatorPointerType data) { |
| typedef std::conditional_t<Self::PacketAccess, typename Self::PacketReturnType, CoeffReturnType> OutType; |
| static_assert(!((EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1) & |
| (EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1 - 1)), |
| "The Local thread size must be a power of 2 for the reduction " |
| "operation"); |
| EIGEN_CONSTEXPR Index local_range = EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1; |
| |
| typename Self::Index inputSize = self.impl().dimensions().TotalSize(); |
| // 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 512 |
| // elemnts individually, we get better performance. |
| const Index reductionPerThread = 2048; |
| // const Index num_work_group = |
| Index reductionGroup = dev.getPowerOfTwo( |
| (inputSize + (reductionPerThread * local_range - 1)) / (reductionPerThread * local_range), true); |
| const Index num_work_group = std::min(reductionGroup, local_range); |
| // 1 |
| // ? local_range |
| // : 1); |
| const Index global_range = num_work_group * local_range; |
| |
| auto thread_range = cl::sycl::nd_range<1>(cl::sycl::range<1>(global_range), cl::sycl::range<1>(local_range)); |
| typedef TensorSycl::internal::FullReductionKernelFunctor<Self, Op, local_range> reduction_kernel_t; |
| if (num_work_group > 1) { |
| CoeffReturnType *temp_pointer = |
| static_cast<CoeffReturnType *>(dev.allocate_temp(num_work_group * sizeof(CoeffReturnType))); |
| typename Self::EvaluatorPointerType tmp_global_accessor = dev.get(temp_pointer); |
| dev.template unary_kernel_launcher<OutType, reduction_kernel_t>(self, tmp_global_accessor, thread_range, |
| local_range, inputSize, reducer) |
| .wait(); |
| typedef TensorSycl::internal::SecondStepFullReducer<CoeffReturnType, Op, EvaluatorPointerType, |
| EvaluatorPointerType, Index, local_range> |
| GenericRKernel; |
| dev.template unary_kernel_launcher<CoeffReturnType, GenericRKernel>( |
| tmp_global_accessor, data, |
| cl::sycl::nd_range<1>(cl::sycl::range<1>(num_work_group), cl::sycl::range<1>(num_work_group)), |
| num_work_group, reducer) |
| .wait(); |
| dev.deallocate_temp(temp_pointer); |
| } else { |
| dev.template unary_kernel_launcher<OutType, reduction_kernel_t>(self, data, thread_range, local_range, inputSize, |
| reducer) |
| .wait(); |
| } |
| } |
| }; |
| // vectorizable inner_most most dim preserver |
| // col reduction |
| template <typename Self, typename Op> |
| struct OuterReducer<Self, Op, Eigen::SyclDevice> { |
| static EIGEN_CONSTEXPR bool HasOptimizedImplementation = true; |
| |
| static bool run(const Self &self, const Op &reducer, const Eigen::SyclDevice &dev, |
| typename Self::EvaluatorPointerType output, typename Self::Index num_coeffs_to_reduce, |
| typename Self::Index num_coeffs_to_preserve) { |
| return ::Eigen::TensorSycl::internal::PartialReducerLauncher< |
| Self, Op, ::Eigen::TensorSycl::internal::reduction_dim::outer_most>::run(self, reducer, dev, output, |
| num_coeffs_to_reduce, |
| num_coeffs_to_preserve); |
| } |
| }; |
| // row reduction |
| template <typename Self, typename Op> |
| struct InnerReducer<Self, Op, Eigen::SyclDevice> { |
| static EIGEN_CONSTEXPR bool HasOptimizedImplementation = true; |
| |
| static bool run(const Self &self, const Op &reducer, const Eigen::SyclDevice &dev, |
| typename Self::EvaluatorPointerType output, typename Self::Index num_coeffs_to_reduce, |
| typename Self::Index num_coeffs_to_preserve) { |
| return ::Eigen::TensorSycl::internal::PartialReducerLauncher< |
| Self, Op, ::Eigen::TensorSycl::internal::reduction_dim::inner_most>::run(self, reducer, dev, output, |
| num_coeffs_to_reduce, |
| num_coeffs_to_preserve); |
| } |
| }; |
| |
| // ArmgMax uses this kernel for partial reduction// |
| // TODO(@mehdi.goli) come up with a better kernel |
| // generic partial reduction |
| template <typename Self, typename Op> |
| struct GenericReducer<Self, Op, Eigen::SyclDevice> { |
| static EIGEN_CONSTEXPR bool HasOptimizedImplementation = false; |
| static bool run(const Self &self, const Op &reducer, const Eigen::SyclDevice &dev, |
| typename Self::EvaluatorPointerType output, typename Self::Index num_values_to_reduce, |
| typename Self::Index num_coeffs_to_preserve) { |
| typename Self::Index range, GRange, tileSize; |
| dev.parallel_for_setup(num_coeffs_to_preserve, tileSize, range, GRange); |
| |
| dev.template unary_kernel_launcher<typename Self::CoeffReturnType, |
| TensorSycl::internal::GenericNondeterministicReducer<Self, Op>>( |
| self, output, cl::sycl::nd_range<1>(cl::sycl::range<1>(GRange), cl::sycl::range<1>(tileSize)), Index(1), |
| reducer, range, (num_values_to_reduce != 0) ? num_values_to_reduce : static_cast<Index>(1)) |
| .wait(); |
| return false; |
| } |
| }; |
| |
| } // namespace internal |
| } // namespace Eigen |
| |
| #endif // UNSUPPORTED_EIGEN_CXX11_SRC_TENSOR_TENSOR_REDUCTION_SYCL_HPP |
