11 #ifndef EIGEN_CXX11_TENSOR_TENSOR_REDUCTION_H 12 #define EIGEN_CXX11_TENSOR_TENSOR_REDUCTION_H 24 template<
typename Op,
typename Dims,
typename XprType,
template <
class>
class MakePointer_ >
25 struct traits<TensorReductionOp<Op, Dims, XprType, MakePointer_> >
28 typedef traits<XprType> XprTraits;
29 typedef typename XprTraits::Scalar Scalar;
30 typedef typename XprTraits::StorageKind StorageKind;
31 typedef typename XprTraits::Index Index;
32 typedef typename XprType::Nested Nested;
33 static const int NumDimensions = XprTraits::NumDimensions - array_size<Dims>::value;
34 static const int Layout = XprTraits::Layout;
36 template <
class T>
struct MakePointer {
38 typedef MakePointer_<T> MakePointerT;
39 typedef typename MakePointerT::Type Type;
43 template<
typename Op,
typename Dims,
typename XprType,
template <
class>
class MakePointer_>
44 struct eval<TensorReductionOp<Op, Dims, XprType, MakePointer_>,
Eigen::Dense>
46 typedef const TensorReductionOp<Op, Dims, XprType, MakePointer_>& type;
49 template<
typename Op,
typename Dims,
typename XprType,
template <
class>
class MakePointer_>
50 struct nested<TensorReductionOp<Op, Dims, XprType, MakePointer_>, 1, typename eval<TensorReductionOp<Op, Dims, XprType, MakePointer_> >::type>
52 typedef TensorReductionOp<Op, Dims, XprType, MakePointer_> type;
56 template <
typename OutputDims>
struct DimInitializer {
57 template <
typename InputDims,
typename ReducedDims> EIGEN_DEVICE_FUNC
58 static void run(
const InputDims& input_dims,
59 const array<
bool, internal::array_size<InputDims>::value>& reduced,
60 OutputDims* output_dims, ReducedDims* reduced_dims) {
61 const int NumInputDims = internal::array_size<InputDims>::value;
64 for (
int i = 0; i < NumInputDims; ++i) {
66 (*reduced_dims)[reduceIndex] = input_dims[i];
69 (*output_dims)[outputIndex] = input_dims[i];
76 template <>
struct DimInitializer<Sizes<> > {
77 template <
typename InputDims,
typename Index,
size_t Rank> EIGEN_DEVICE_FUNC
78 static void run(
const InputDims& input_dims,
const array<bool, Rank>&,
79 Sizes<>*, array<Index, Rank>* reduced_dims) {
80 const int NumInputDims = internal::array_size<InputDims>::value;
81 for (
int i = 0; i < NumInputDims; ++i) {
82 (*reduced_dims)[i] = input_dims[i];
88 template <
typename ReducedDims,
int NumTensorDims,
int Layout>
89 struct are_inner_most_dims {
90 static const bool value =
false;
92 template <
typename ReducedDims,
int NumTensorDims,
int Layout>
93 struct preserve_inner_most_dims {
94 static const bool value =
false;
97 #if EIGEN_HAS_CONSTEXPR && EIGEN_HAS_VARIADIC_TEMPLATES 98 template <
typename ReducedDims,
int NumTensorDims>
99 struct are_inner_most_dims<ReducedDims, NumTensorDims, ColMajor>{
100 static const bool tmp1 = indices_statically_known_to_increase<ReducedDims>();
101 static const bool tmp2 = index_statically_eq<ReducedDims>(0, 0);
102 static const bool tmp3 = index_statically_eq<ReducedDims>(array_size<ReducedDims>::value-1, array_size<ReducedDims>::value-1);
103 static const bool value = tmp1 & tmp2 & tmp3;
105 template <
typename ReducedDims,
int NumTensorDims>
106 struct are_inner_most_dims<ReducedDims, NumTensorDims, RowMajor>{
107 static const bool tmp1 = indices_statically_known_to_increase<ReducedDims>();
108 static const bool tmp2 = index_statically_eq<ReducedDims>(0, NumTensorDims - array_size<ReducedDims>::value);
109 static const bool tmp3 = index_statically_eq<ReducedDims>(array_size<ReducedDims>::value - 1, NumTensorDims - 1);
110 static const bool value = tmp1 & tmp2 & tmp3;
113 template <
typename ReducedDims,
int NumTensorDims>
114 struct preserve_inner_most_dims<ReducedDims, NumTensorDims, ColMajor>{
115 static const bool tmp1 = indices_statically_known_to_increase<ReducedDims>();
116 static const bool tmp2 = index_statically_gt<ReducedDims>(0, 0);
117 static const bool value = tmp1 & tmp2;
120 template <
typename ReducedDims,
int NumTensorDims>
121 struct preserve_inner_most_dims<ReducedDims, NumTensorDims, RowMajor>{
122 static const bool tmp1 = indices_statically_known_to_increase<ReducedDims>();
123 static const bool tmp2 = index_statically_lt<ReducedDims>(array_size<ReducedDims>::value - 1, NumTensorDims - 1);
124 static const bool value = tmp1 & tmp2;
129 template <
int DimIndex,
typename Self,
typename Op>
130 struct GenericDimReducer {
131 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&
self,
typename Self::Index firstIndex, Op& reducer,
typename Self::CoeffReturnType* accum) {
132 EIGEN_STATIC_ASSERT((DimIndex > 0), YOU_MADE_A_PROGRAMMING_MISTAKE);
133 for (
int j = 0; j <
self.m_reducedDims[DimIndex]; ++j) {
134 const typename Self::Index input = firstIndex + j *
self.m_reducedStrides[DimIndex];
135 GenericDimReducer<DimIndex-1, Self, Op>::reduce(
self, input, reducer, accum);
139 template <
typename Self,
typename Op>
140 struct GenericDimReducer<0, Self, Op> {
141 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&
self,
typename Self::Index firstIndex, Op& reducer,
typename Self::CoeffReturnType* accum) {
142 for (
int j = 0; j <
self.m_reducedDims[0]; ++j) {
143 const typename Self::Index input = firstIndex + j *
self.m_reducedStrides[0];
144 reducer.reduce(
self.m_impl.coeff(input), accum);
148 template <
typename Self,
typename Op>
149 struct GenericDimReducer<-1, Self, Op> {
150 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&
self,
typename Self::Index index, Op& reducer,
typename Self::CoeffReturnType* accum) {
151 reducer.reduce(
self.m_impl.coeff(index), accum);
155 template <
typename Self,
typename Op,
bool Vectorizable = (Self::InputPacketAccess & Op::PacketAccess)>
156 struct InnerMostDimReducer {
157 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
typename Self::CoeffReturnType reduce(
const Self&
self,
typename Self::Index firstIndex,
typename Self::Index numValuesToReduce, Op& reducer) {
158 typename Self::CoeffReturnType accum = reducer.initialize();
159 for (
typename Self::Index j = 0; j < numValuesToReduce; ++j) {
160 reducer.reduce(
self.m_impl.coeff(firstIndex + j), &accum);
162 return reducer.finalize(accum);
166 template <
typename Self,
typename Op>
167 struct InnerMostDimReducer<Self, Op, true> {
168 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
typename Self::CoeffReturnType reduce(
const Self&
self,
typename Self::Index firstIndex,
typename Self::Index numValuesToReduce, Op& reducer) {
169 const int packetSize = internal::unpacket_traits<typename Self::PacketReturnType>::size;
170 const typename Self::Index VectorizedSize = (numValuesToReduce / packetSize) * packetSize;
171 typename Self::PacketReturnType p = reducer.template initializePacket<typename Self::PacketReturnType>();
172 for (
typename Self::Index j = 0; j < VectorizedSize; j += packetSize) {
173 reducer.reducePacket(
self.m_impl.template packet<Unaligned>(firstIndex + j), &p);
175 typename Self::CoeffReturnType accum = reducer.initialize();
176 for (
typename Self::Index j = VectorizedSize; j < numValuesToReduce; ++j) {
177 reducer.reduce(
self.m_impl.coeff(firstIndex + j), &accum);
179 return reducer.finalizeBoth(accum, p);
183 template <
int DimIndex,
typename Self,
typename Op,
bool vectorizable = (Self::InputPacketAccess & Op::PacketAccess)>
184 struct InnerMostDimPreserver {
185 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&,
typename Self::Index, Op&,
typename Self::PacketReturnType*) {
186 eigen_assert(
false &&
"should never be called");
190 template <
int DimIndex,
typename Self,
typename Op>
191 struct InnerMostDimPreserver<DimIndex, Self, Op, true> {
192 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&
self,
typename Self::Index firstIndex, Op& reducer,
typename Self::PacketReturnType* accum) {
193 EIGEN_STATIC_ASSERT((DimIndex > 0), YOU_MADE_A_PROGRAMMING_MISTAKE);
194 for (
typename Self::Index j = 0; j <
self.m_reducedDims[DimIndex]; ++j) {
195 const typename Self::Index input = firstIndex + j *
self.m_reducedStrides[DimIndex];
196 InnerMostDimPreserver<DimIndex-1, Self, Op>::reduce(
self, input, reducer, accum);
201 template <
typename Self,
typename Op>
202 struct InnerMostDimPreserver<0, Self, Op, true> {
203 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&
self,
typename Self::Index firstIndex, Op& reducer,
typename Self::PacketReturnType* accum) {
204 for (
typename Self::Index j = 0; j <
self.m_reducedDims[0]; ++j) {
205 const typename Self::Index input = firstIndex + j *
self.m_reducedStrides[0];
206 reducer.reducePacket(
self.m_impl.template packet<Unaligned>(input), accum);
210 template <
typename Self,
typename Op>
211 struct InnerMostDimPreserver<-1, Self, Op, true> {
212 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&,
typename Self::Index, Op&,
typename Self::PacketReturnType*) {
213 eigen_assert(
false &&
"should never be called");
218 template <
typename Self,
typename Op,
typename Device,
bool Vectorizable = (Self::InputPacketAccess & Op::PacketAccess)>
220 static const bool HasOptimizedImplementation =
false;
222 static EIGEN_DEVICE_FUNC
void run(
const Self&
self, Op& reducer,
const Device&,
typename Self::CoeffReturnType* output) {
223 const typename Self::Index num_coeffs = array_prod(
self.m_impl.dimensions());
224 *output = InnerMostDimReducer<Self, Op, Vectorizable>::reduce(
self, 0, num_coeffs, reducer);
229 #ifdef EIGEN_USE_THREADS 231 template <
typename Self,
typename Op,
232 bool Vectorizable = (Self::InputPacketAccess & Op::PacketAccess)>
233 struct FullReducerShard {
234 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void run(
const Self&
self,
typename Self::Index firstIndex,
235 typename Self::Index numValuesToReduce, Op& reducer,
236 typename Self::CoeffReturnType* output) {
237 *output = InnerMostDimReducer<Self, Op, Vectorizable>::reduce(
238 self, firstIndex, numValuesToReduce, reducer);
243 template <
typename Self,
typename Op,
bool Vectorizable>
244 struct FullReducer<Self, Op, ThreadPoolDevice, Vectorizable> {
245 static const bool HasOptimizedImplementation = !Op::IsStateful;
246 static const int PacketSize =
247 unpacket_traits<typename Self::PacketReturnType>::size;
250 static void run(
const Self&
self, Op& reducer,
const ThreadPoolDevice& device,
251 typename Self::CoeffReturnType* output) {
252 typedef typename Self::Index Index;
253 const Index num_coeffs = array_prod(
self.m_impl.dimensions());
254 if (num_coeffs == 0) {
255 *output = reducer.finalize(reducer.initialize());
258 const TensorOpCost cost =
259 self.m_impl.costPerCoeff(Vectorizable) +
260 TensorOpCost(0, 0, internal::functor_traits<Op>::Cost, Vectorizable,
262 const int num_threads = TensorCostModel<ThreadPoolDevice>::numThreads(
263 num_coeffs, cost, device.numThreads());
264 if (num_threads == 1) {
266 InnerMostDimReducer<Self, Op, Vectorizable>::reduce(
self, 0, num_coeffs, reducer);
269 const Index blocksize =
270 std::floor<Index>(
static_cast<float>(num_coeffs) / num_threads);
271 const Index numblocks = blocksize > 0 ? num_coeffs / blocksize : 0;
272 eigen_assert(num_coeffs >= numblocks * blocksize);
274 Barrier barrier(internal::convert_index<unsigned int>(numblocks));
275 MaxSizeVector<typename Self::CoeffReturnType> shards(numblocks, reducer.initialize());
276 for (Index i = 0; i < numblocks; ++i) {
277 device.enqueue_with_barrier(&barrier, &FullReducerShard<Self, Op, Vectorizable>::run,
278 self, i * blocksize, blocksize, reducer,
281 typename Self::CoeffReturnType finalShard;
282 if (numblocks * blocksize < num_coeffs) {
283 finalShard = InnerMostDimReducer<Self, Op, Vectorizable>::reduce(
284 self, numblocks * blocksize, num_coeffs - numblocks * blocksize,
287 finalShard = reducer.initialize();
291 for (Index i = 0; i < numblocks; ++i) {
292 reducer.reduce(shards[i], &finalShard);
294 *output = reducer.finalize(finalShard);
302 template <
typename Self,
typename Op,
typename Device>
303 struct InnerReducer {
304 static const bool HasOptimizedImplementation =
false;
306 EIGEN_DEVICE_FUNC
static bool run(
const Self&, Op&,
const Device&,
typename Self::CoeffReturnType*,
typename Self::Index,
typename Self::Index) {
307 eigen_assert(
false &&
"Not implemented");
313 template <
typename Self,
typename Op,
typename Device>
314 struct OuterReducer {
315 static const bool HasOptimizedImplementation =
false;
317 EIGEN_DEVICE_FUNC
static bool run(
const Self&, Op&,
const Device&,
typename Self::CoeffReturnType*,
typename Self::Index,
typename Self::Index) {
318 eigen_assert(
false &&
"Not implemented");
324 #if defined(EIGEN_USE_GPU) && defined(__CUDACC__) 325 template <
int B,
int N,
typename S,
typename R,
typename I>
326 __global__
void FullReductionKernel(R,
const S, I,
typename S::CoeffReturnType*,
unsigned int*);
329 #ifdef EIGEN_HAS_CUDA_FP16 330 template <
typename S,
typename R,
typename I>
331 __global__
void ReductionInitFullReduxKernelHalfFloat(R,
const S, I, half2*);
332 template <
int B,
int N,
typename S,
typename R,
typename I>
333 __global__
void FullReductionKernelHalfFloat(R,
const S, I, half*, half2*);
334 template <
int NPT,
typename S,
typename R,
typename I>
335 __global__
void InnerReductionKernelHalfFloat(R,
const S, I, I, half*);
339 template <
int NPT,
typename S,
typename R,
typename I>
340 __global__
void InnerReductionKernel(R,
const S, I, I,
typename S::CoeffReturnType*);
342 template <
int NPT,
typename S,
typename R,
typename I>
343 __global__
void OuterReductionKernel(R,
const S, I, I,
typename S::CoeffReturnType*);
349 template <
typename Op,
typename Dims,
typename XprType,
template <
class>
class MakePointer_>
350 class TensorReductionOp :
public TensorBase<TensorReductionOp<Op, Dims, XprType, MakePointer_>, ReadOnlyAccessors> {
352 typedef typename Eigen::internal::traits<TensorReductionOp>::Scalar Scalar;
353 typedef typename Eigen::NumTraits<Scalar>::Real RealScalar;
354 typedef typename internal::remove_const<typename XprType::CoeffReturnType>::type CoeffReturnType;
355 typedef typename Eigen::internal::nested<TensorReductionOp>::type Nested;
356 typedef typename Eigen::internal::traits<TensorReductionOp>::StorageKind StorageKind;
357 typedef typename Eigen::internal::traits<TensorReductionOp>::Index Index;
359 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
360 TensorReductionOp(
const XprType& expr,
const Dims& dims) : m_expr(expr), m_dims(dims)
362 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
363 TensorReductionOp(
const XprType& expr,
const Dims& dims,
const Op& reducer) : m_expr(expr), m_dims(dims), m_reducer(reducer)
366 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
367 const XprType& expression()
const {
return m_expr; }
368 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
369 const Dims& dims()
const {
return m_dims; }
370 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
371 const Op& reducer()
const {
return m_reducer; }
374 typename XprType::Nested m_expr;
381 template<
typename Op,
typename Dims,
typename ArgType,
template <
class>
class MakePointer_,
typename Device>
382 struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType, MakePointer_>, Device>
384 typedef TensorReductionOp<Op, Dims, ArgType, MakePointer_> XprType;
385 typedef typename XprType::Index Index;
386 typedef ArgType ChildType;
387 typedef typename TensorEvaluator<ArgType, Device>::Dimensions InputDimensions;
388 static const int NumInputDims = internal::array_size<InputDimensions>::value;
389 static const int NumReducedDims = internal::array_size<Dims>::value;
390 static const int NumOutputDims = NumInputDims - NumReducedDims;
391 typedef typename internal::conditional<NumOutputDims==0, Sizes<>, DSizes<Index, NumOutputDims> >::type Dimensions;
392 typedef typename XprType::Scalar Scalar;
393 typedef TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType, MakePointer_>, Device> Self;
394 static const bool InputPacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess;
395 typedef typename internal::remove_const<typename XprType::CoeffReturnType>::type CoeffReturnType;
396 typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
397 static const int PacketSize = internal::unpacket_traits<PacketReturnType>::size;
401 PacketAccess = Self::InputPacketAccess && Op::PacketAccess,
402 Layout = TensorEvaluator<ArgType, Device>::Layout,
407 static const bool ReducingInnerMostDims = internal::are_inner_most_dims<Dims, NumInputDims, Layout>::value;
408 static const bool PreservingInnerMostDims = internal::preserve_inner_most_dims<Dims, NumInputDims, Layout>::value;
409 static const bool RunningFullReduction = (NumOutputDims==0);
411 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator(
const XprType& op,
const Device& device)
412 : m_impl(op.expression(), device), m_reducer(op.reducer()), m_result(NULL), m_device(device), m_xpr_dims(op.dims())
414 EIGEN_STATIC_ASSERT((NumInputDims >= NumReducedDims), YOU_MADE_A_PROGRAMMING_MISTAKE);
415 EIGEN_STATIC_ASSERT((!ReducingInnerMostDims | !PreservingInnerMostDims | (NumReducedDims == NumInputDims)),
416 YOU_MADE_A_PROGRAMMING_MISTAKE);
419 for (
int i = 0; i < NumInputDims; ++i) {
420 m_reduced[i] =
false;
422 for (
int i = 0; i < NumReducedDims; ++i) {
423 eigen_assert(op.dims()[i] >= 0);
424 eigen_assert(op.dims()[i] < NumInputDims);
425 m_reduced[op.dims()[i]] =
true;
428 const typename TensorEvaluator<ArgType, Device>::Dimensions& input_dims = m_impl.dimensions();
429 internal::DimInitializer<Dimensions>::run(input_dims, m_reduced, &m_dimensions, &m_reducedDims);
432 if (NumOutputDims > 0) {
433 if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
434 m_outputStrides[0] = 1;
435 for (
int i = 1; i < NumOutputDims; ++i) {
436 m_outputStrides[i] = m_outputStrides[i - 1] * m_dimensions[i - 1];
439 m_outputStrides.back() = 1;
440 for (
int i = NumOutputDims - 2; i >= 0; --i) {
441 m_outputStrides[i] = m_outputStrides[i + 1] * m_dimensions[i + 1];
447 if (NumInputDims > 0) {
448 array<Index, NumInputDims> input_strides;
449 if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
450 input_strides[0] = 1;
451 for (
int i = 1; i < NumInputDims; ++i) {
452 input_strides[i] = input_strides[i-1] * input_dims[i-1];
455 input_strides.back() = 1;
456 for (
int i = NumInputDims - 2; i >= 0; --i) {
457 input_strides[i] = input_strides[i + 1] * input_dims[i + 1];
463 for (
int i = 0; i < NumInputDims; ++i) {
465 m_reducedStrides[reduceIndex] = input_strides[i];
468 m_preservedStrides[outputIndex] = input_strides[i];
475 if (NumOutputDims == 0) {
476 m_preservedStrides[0] = internal::array_prod(input_dims);
480 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
const Dimensions& dimensions()
const {
return m_dimensions; }
482 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool evalSubExprsIfNeeded(
typename MakePointer_<CoeffReturnType>::Type data) {
483 m_impl.evalSubExprsIfNeeded(NULL);
486 if ((RunningFullReduction && RunningOnSycl) ||(RunningFullReduction &&
487 internal::FullReducer<Self, Op, Device>::HasOptimizedImplementation &&
488 ((RunningOnGPU && (m_device.majorDeviceVersion() >= 3)) ||
490 bool need_assign =
false;
492 m_result =
static_cast<CoeffReturnType*
>(m_device.allocate(
sizeof(CoeffReturnType)));
496 Op reducer(m_reducer);
497 internal::FullReducer<Self, Op, Device>::run(*
this, reducer, m_device, data);
500 else if(RunningOnSycl){
501 const Index num_values_to_reduce = internal::array_prod(m_reducedDims);
502 const Index num_coeffs_to_preserve = internal::array_prod(m_dimensions);
504 data =
static_cast<CoeffReturnType*
>(m_device.allocate(
sizeof(CoeffReturnType) * num_coeffs_to_preserve));
507 Op reducer(m_reducer);
508 internal::InnerReducer<Self, Op, Device>::run(*
this, reducer, m_device, data, num_values_to_reduce, num_coeffs_to_preserve);
509 return (m_result != NULL);
513 else if (RunningOnGPU && (m_device.majorDeviceVersion() >= 3)) {
514 bool reducing_inner_dims =
true;
515 for (
int i = 0; i < NumReducedDims; ++i) {
516 if (static_cast<int>(Layout) ==
static_cast<int>(ColMajor)) {
517 reducing_inner_dims &= m_reduced[i];
519 reducing_inner_dims &= m_reduced[NumInputDims - 1 - i];
522 if (internal::InnerReducer<Self, Op, Device>::HasOptimizedImplementation &&
523 (reducing_inner_dims || ReducingInnerMostDims)) {
524 const Index num_values_to_reduce = internal::array_prod(m_reducedDims);
525 const Index num_coeffs_to_preserve = internal::array_prod(m_dimensions);
527 if (num_coeffs_to_preserve < 1024 && num_values_to_reduce > num_coeffs_to_preserve && num_values_to_reduce > 128) {
528 data =
static_cast<CoeffReturnType*
>(m_device.allocate(
sizeof(CoeffReturnType) * num_coeffs_to_preserve));
535 Op reducer(m_reducer);
536 if (internal::InnerReducer<Self, Op, Device>::run(*
this, reducer, m_device, data, num_values_to_reduce, num_coeffs_to_preserve)) {
538 m_device.deallocate(m_result);
543 return (m_result != NULL);
547 bool preserving_inner_dims =
true;
548 for (
int i = 0; i < NumReducedDims; ++i) {
549 if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
550 preserving_inner_dims &= m_reduced[NumInputDims - 1 - i];
552 preserving_inner_dims &= m_reduced[i];
555 if (internal::OuterReducer<Self, Op, Device>::HasOptimizedImplementation &&
556 preserving_inner_dims) {
557 const Index num_values_to_reduce = internal::array_prod(m_reducedDims);
558 const Index num_coeffs_to_preserve = internal::array_prod(m_dimensions);
560 if (num_coeffs_to_preserve < 1024 && num_values_to_reduce > num_coeffs_to_preserve && num_values_to_reduce > 32) {
561 data =
static_cast<CoeffReturnType*
>(m_device.allocate(
sizeof(CoeffReturnType) * num_coeffs_to_preserve));
568 Op reducer(m_reducer);
569 if (internal::OuterReducer<Self, Op, Device>::run(*
this, reducer, m_device, data, num_values_to_reduce, num_coeffs_to_preserve)) {
571 m_device.deallocate(m_result);
576 return (m_result != NULL);
583 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void cleanup() {
586 m_device.deallocate(m_result);
591 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index)
const 593 if ((RunningOnSycl || RunningFullReduction || RunningOnGPU) && m_result) {
594 return *(m_result + index);
596 Op reducer(m_reducer);
597 if (ReducingInnerMostDims || RunningFullReduction) {
598 const Index num_values_to_reduce =
599 (
static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1];
600 return internal::InnerMostDimReducer<Self, Op>::reduce(*
this, firstInput(index),
601 num_values_to_reduce, reducer);
603 typename Self::CoeffReturnType accum = reducer.initialize();
604 internal::GenericDimReducer<NumReducedDims-1, Self, Op>::reduce(*
this, firstInput(index), reducer, &accum);
605 return reducer.finalize(accum);
610 template<
int LoadMode>
611 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType packet(Index index)
const 613 EIGEN_STATIC_ASSERT((PacketSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE)
614 eigen_assert(index + PacketSize - 1 < Index(internal::array_prod(dimensions())));
616 if (RunningOnGPU && m_result) {
617 return internal::pload<PacketReturnType>(m_result + index);
620 EIGEN_ALIGN_MAX
typename internal::remove_const<CoeffReturnType>::type values[PacketSize];
621 if (ReducingInnerMostDims) {
622 const Index num_values_to_reduce =
623 (
static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1];
624 const Index firstIndex = firstInput(index);
625 for (Index i = 0; i < PacketSize; ++i) {
626 Op reducer(m_reducer);
627 values[i] = internal::InnerMostDimReducer<Self, Op>::reduce(*
this, firstIndex + i * num_values_to_reduce,
628 num_values_to_reduce, reducer);
630 }
else if (PreservingInnerMostDims) {
631 const Index firstIndex = firstInput(index);
632 const int innermost_dim = (
static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? 0 : NumOutputDims - 1;
634 if (((firstIndex % m_dimensions[innermost_dim]) + PacketSize - 1) < m_dimensions[innermost_dim]) {
635 Op reducer(m_reducer);
636 typename Self::PacketReturnType accum = reducer.template initializePacket<typename Self::PacketReturnType>();
637 internal::InnerMostDimPreserver<NumReducedDims-1, Self, Op>::reduce(*
this, firstIndex, reducer, &accum);
638 return reducer.finalizePacket(accum);
640 for (
int i = 0; i < PacketSize; ++i) {
641 values[i] = coeff(index + i);
645 for (
int i = 0; i < PacketSize; ++i) {
646 values[i] = coeff(index + i);
649 PacketReturnType rslt = internal::pload<PacketReturnType>(values);
654 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(
bool vectorized)
const {
655 if (RunningFullReduction && m_result) {
656 return TensorOpCost(
sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize);
658 const Index num_values_to_reduce = internal::array_prod(m_reducedDims);
659 const double compute_cost = num_values_to_reduce * internal::functor_traits<Op>::Cost;
660 return m_impl.costPerCoeff(vectorized) * num_values_to_reduce +
661 TensorOpCost(0, 0, compute_cost, vectorized, PacketSize);
665 EIGEN_DEVICE_FUNC
typename MakePointer_<Scalar>::Type data()
const {
return m_result; }
667 const TensorEvaluator<ArgType, Device>& impl()
const {
return m_impl; }
669 const Device& device()
const{
return m_device;}
671 const Dims& xprDims()
const {
return m_xpr_dims;}
675 template <
int,
typename,
typename>
friend struct internal::GenericDimReducer;
676 template <
typename,
typename,
bool>
friend struct internal::InnerMostDimReducer;
677 template <
int,
typename,
typename,
bool>
friend struct internal::InnerMostDimPreserver;
678 template <
typename S,
typename O,
typename D,
bool V>
friend struct internal::FullReducer;
679 #ifdef EIGEN_USE_THREADS 680 template <
typename S,
typename O,
bool V>
friend struct internal::FullReducerShard;
682 #if defined(EIGEN_USE_GPU) && defined(__CUDACC__) 683 template <
int B,
int N,
typename S,
typename R,
typename I>
friend void internal::FullReductionKernel(R,
const S, I,
typename S::CoeffReturnType*,
unsigned int*);
684 #ifdef EIGEN_HAS_CUDA_FP16 685 template <
typename S,
typename R,
typename I>
friend void internal::ReductionInitFullReduxKernelHalfFloat(R,
const S, I, half2*);
686 template <
int B,
int N,
typename S,
typename R,
typename I>
friend void internal::FullReductionKernelHalfFloat(R,
const S, I, half*, half2*);
687 template <
int NPT,
typename S,
typename R,
typename I>
friend void internal::InnerReductionKernelHalfFloat(R,
const S, I, I, half*);
689 template <
int NPT,
typename S,
typename R,
typename I>
friend void internal::InnerReductionKernel(R,
const S, I, I,
typename S::CoeffReturnType*);
691 template <
int NPT,
typename S,
typename R,
typename I>
friend void internal::OuterReductionKernel(R,
const S, I, I,
typename S::CoeffReturnType*);
694 template <
typename S,
typename O,
typename D>
friend struct internal::InnerReducer;
698 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index firstInput(Index index)
const {
699 if (ReducingInnerMostDims) {
700 if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
701 return index * m_preservedStrides[0];
703 return index * m_preservedStrides[NumPreservedStrides - 1];
707 Index startInput = 0;
708 if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
709 for (
int i = NumOutputDims - 1; i > 0; --i) {
711 const Index idx = index / m_outputStrides[i];
712 startInput += idx * m_preservedStrides[i];
713 index -= idx * m_outputStrides[i];
715 if (PreservingInnerMostDims) {
716 eigen_assert(m_preservedStrides[0] == 1);
719 startInput += index * m_preservedStrides[0];
722 for (
int i = 0; i < NumOutputDims - 1; ++i) {
724 const Index idx = index / m_outputStrides[i];
725 startInput += idx * m_preservedStrides[i];
726 index -= idx * m_outputStrides[i];
728 if (PreservingInnerMostDims) {
729 eigen_assert(m_preservedStrides[NumPreservedStrides - 1] == 1);
732 startInput += index * m_preservedStrides[NumPreservedStrides - 1];
739 array<bool, NumInputDims> m_reduced;
741 Dimensions m_dimensions;
743 array<Index, NumOutputDims> m_outputStrides;
746 static const int NumPreservedStrides = max_n_1<NumOutputDims>::size;
747 array<Index, NumPreservedStrides> m_preservedStrides;
751 array<Index, NumReducedDims> m_reducedStrides;
754 array<Index, NumReducedDims> m_reducedDims;
757 TensorEvaluator<ArgType, Device> m_impl;
763 #if defined(EIGEN_USE_GPU) && defined(__CUDACC__) 764 static const bool RunningOnGPU = internal::is_same<Device, Eigen::GpuDevice>::value;
765 static const bool RunningOnSycl =
false;
766 #elif defined(EIGEN_USE_SYCL) 767 static const bool RunningOnSycl = internal::is_same<typename internal::remove_all<Device>::type, Eigen::SyclDevice>::value;
768 static const bool RunningOnGPU =
false;
770 static const bool RunningOnGPU =
false;
771 static const bool RunningOnSycl =
false;
773 typename MakePointer_<CoeffReturnType>::Type m_result;
775 const Device& m_device;
776 const Dims& m_xpr_dims;
781 #endif // EIGEN_CXX11_TENSOR_TENSOR_REDUCTION_H Namespace containing all symbols from the Eigen library.
Definition: AdolcForward:45