10 #ifndef EIGEN_CXX11_TENSOR_TENSOR_REDUCTION_H 11 #define EIGEN_CXX11_TENSOR_TENSOR_REDUCTION_H 23 template<
typename Op,
typename Dims,
typename XprType>
24 struct traits<TensorReductionOp<Op, Dims, XprType> >
27 typedef traits<XprType> XprTraits;
28 typedef typename XprTraits::Scalar Scalar;
29 typedef typename XprTraits::StorageKind StorageKind;
30 typedef typename XprTraits::Index Index;
31 typedef typename XprType::Nested Nested;
32 static const int NumDimensions = XprTraits::NumDimensions - array_size<Dims>::value;
33 static const int Layout = XprTraits::Layout;
36 template<
typename Op,
typename Dims,
typename XprType>
37 struct eval<TensorReductionOp<Op, Dims, XprType>,
Eigen::Dense>
39 typedef const TensorReductionOp<Op, Dims, XprType>& type;
42 template<
typename Op,
typename Dims,
typename XprType>
43 struct nested<TensorReductionOp<Op, Dims, XprType>, 1, typename eval<TensorReductionOp<Op, Dims, XprType> >::type>
45 typedef TensorReductionOp<Op, Dims, XprType> type;
49 template <
typename OutputDims>
struct DimInitializer {
50 template <
typename InputDims,
typename ReducedDims> EIGEN_DEVICE_FUNC
51 static void run(
const InputDims& input_dims,
52 const array<
bool, internal::array_size<InputDims>::value>& reduced,
53 OutputDims* output_dims, ReducedDims* reduced_dims) {
54 const int NumInputDims = internal::array_size<InputDims>::value;
57 for (
int i = 0; i < NumInputDims; ++i) {
59 (*reduced_dims)[reduceIndex] = input_dims[i];
62 (*output_dims)[outputIndex] = input_dims[i];
69 template <>
struct DimInitializer<Sizes<> > {
70 template <
typename InputDims,
typename Index,
size_t Rank> EIGEN_DEVICE_FUNC
71 static void run(
const InputDims& input_dims,
const array<bool, Rank>&,
72 Sizes<>*, array<Index, Rank>* reduced_dims) {
73 const int NumInputDims = internal::array_size<InputDims>::value;
74 for (
int i = 0; i < NumInputDims; ++i) {
75 (*reduced_dims)[i] = input_dims[i];
81 template <
typename ReducedDims,
int NumTensorDims,
int Layout>
82 struct are_inner_most_dims {
83 static const bool value =
false;
85 template <
typename ReducedDims,
int NumTensorDims,
int Layout>
86 struct preserve_inner_most_dims {
87 static const bool value =
false;
90 #if EIGEN_HAS_CONSTEXPR && EIGEN_HAS_VARIADIC_TEMPLATES 91 template <
typename ReducedDims,
int NumTensorDims>
92 struct are_inner_most_dims<ReducedDims, NumTensorDims, ColMajor>{
93 static const bool tmp1 = indices_statically_known_to_increase<ReducedDims>();
94 static const bool tmp2 = index_statically_eq<ReducedDims>(0, 0);
95 static const bool tmp3 = index_statically_eq<ReducedDims>(array_size<ReducedDims>::value-1, array_size<ReducedDims>::value-1);
96 static const bool value = tmp1 & tmp2 & tmp3;
98 template <
typename ReducedDims,
int NumTensorDims>
99 struct are_inner_most_dims<ReducedDims, NumTensorDims, RowMajor>{
100 static const bool tmp1 = indices_statically_known_to_increase<ReducedDims>();
101 static const bool tmp2 = index_statically_eq<ReducedDims>(0, NumTensorDims - array_size<ReducedDims>::value);
102 static const bool tmp3 = index_statically_eq<ReducedDims>(array_size<ReducedDims>::value - 1, NumTensorDims - 1);
103 static const bool value = tmp1 & tmp2 & tmp3;
106 template <
typename ReducedDims,
int NumTensorDims>
107 struct preserve_inner_most_dims<ReducedDims, NumTensorDims, ColMajor>{
108 static const bool tmp1 = indices_statically_known_to_increase<ReducedDims>();
109 static const bool tmp2 = index_statically_gt<ReducedDims>(0, 0);
110 static const bool value = tmp1 & tmp2;
113 template <
typename ReducedDims,
int NumTensorDims>
114 struct preserve_inner_most_dims<ReducedDims, NumTensorDims, RowMajor>{
115 static const bool tmp1 = indices_statically_known_to_increase<ReducedDims>();
116 static const bool tmp2 = index_statically_lt<ReducedDims>(array_size<ReducedDims>::value - 1, NumTensorDims - 1);
117 static const bool value = tmp1 & tmp2;
122 template <
int DimIndex,
typename Self,
typename Op>
123 struct GenericDimReducer {
124 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&
self,
typename Self::Index firstIndex, Op& reducer,
typename Self::CoeffReturnType* accum) {
125 EIGEN_STATIC_ASSERT((DimIndex > 0), YOU_MADE_A_PROGRAMMING_MISTAKE);
126 for (
int j = 0; j <
self.m_reducedDims[DimIndex]; ++j) {
127 const typename Self::Index input = firstIndex + j *
self.m_reducedStrides[DimIndex];
128 GenericDimReducer<DimIndex-1, Self, Op>::reduce(
self, input, reducer, accum);
132 template <
typename Self,
typename Op>
133 struct GenericDimReducer<0, Self, Op> {
134 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&
self,
typename Self::Index firstIndex, Op& reducer,
typename Self::CoeffReturnType* accum) {
135 for (
int j = 0; j <
self.m_reducedDims[0]; ++j) {
136 const typename Self::Index input = firstIndex + j *
self.m_reducedStrides[0];
137 reducer.reduce(
self.m_impl.coeff(input), accum);
141 template <
typename Self,
typename Op>
142 struct GenericDimReducer<-1, Self, Op> {
143 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&
self,
typename Self::Index index, Op& reducer,
typename Self::CoeffReturnType* accum) {
144 reducer.reduce(
self.m_impl.coeff(index), accum);
148 template <
typename Self,
typename Op,
bool Vectorizable = (Self::InputPacketAccess & Op::PacketAccess)>
149 struct InnerMostDimReducer {
150 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
typename Self::CoeffReturnType reduce(
const Self&
self,
typename Self::Index firstIndex,
typename Self::Index numValuesToReduce, Op& reducer) {
151 typename Self::CoeffReturnType accum = reducer.initialize();
152 for (
typename Self::Index j = 0; j < numValuesToReduce; ++j) {
153 reducer.reduce(
self.m_impl.coeff(firstIndex + j), &accum);
155 return reducer.finalize(accum);
159 template <
typename Self,
typename Op>
160 struct InnerMostDimReducer<Self, Op, true> {
161 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
typename Self::CoeffReturnType reduce(
const Self&
self,
typename Self::Index firstIndex,
typename Self::Index numValuesToReduce, Op& reducer) {
162 const int packetSize = internal::unpacket_traits<typename Self::PacketReturnType>::size;
163 const typename Self::Index VectorizedSize = (numValuesToReduce / packetSize) * packetSize;
164 typename Self::PacketReturnType p = reducer.template initializePacket<typename Self::PacketReturnType>();
165 for (
typename Self::Index j = 0; j < VectorizedSize; j += packetSize) {
166 reducer.reducePacket(
self.m_impl.template packet<Unaligned>(firstIndex + j), &p);
168 typename Self::CoeffReturnType accum = reducer.initialize();
169 for (
typename Self::Index j = VectorizedSize; j < numValuesToReduce; ++j) {
170 reducer.reduce(
self.m_impl.coeff(firstIndex + j), &accum);
172 return reducer.finalizeBoth(accum, p);
176 template <
int DimIndex,
typename Self,
typename Op,
bool vectorizable = (Self::InputPacketAccess & Op::PacketAccess)>
177 struct InnerMostDimPreserver {
178 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&,
typename Self::Index, Op&,
typename Self::PacketReturnType*) {
179 eigen_assert(
false &&
"should never be called");
183 template <
int DimIndex,
typename Self,
typename Op>
184 struct InnerMostDimPreserver<DimIndex, Self, Op, true> {
185 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&
self,
typename Self::Index firstIndex, Op& reducer,
typename Self::PacketReturnType* accum) {
186 EIGEN_STATIC_ASSERT((DimIndex > 0), YOU_MADE_A_PROGRAMMING_MISTAKE);
187 for (
typename Self::Index j = 0; j <
self.m_reducedDims[DimIndex]; ++j) {
188 const typename Self::Index input = firstIndex + j *
self.m_reducedStrides[DimIndex];
189 InnerMostDimPreserver<DimIndex-1, Self, Op>::reduce(
self, input, reducer, accum);
194 template <
typename Self,
typename Op>
195 struct InnerMostDimPreserver<0, Self, Op, true> {
196 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&
self,
typename Self::Index firstIndex, Op& reducer,
typename Self::PacketReturnType* accum) {
197 for (
typename Self::Index j = 0; j <
self.m_reducedDims[0]; ++j) {
198 const typename Self::Index input = firstIndex + j *
self.m_reducedStrides[0];
199 reducer.reducePacket(
self.m_impl.template packet<Unaligned>(input), accum);
203 template <
typename Self,
typename Op>
204 struct InnerMostDimPreserver<-1, Self, Op, true> {
205 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void reduce(
const Self&,
typename Self::Index, Op&,
typename Self::PacketReturnType*) {
206 eigen_assert(
false &&
"should never be called");
211 template <
typename Self,
typename Op,
typename Device,
bool Vectorizable = (Self::InputPacketAccess & Op::PacketAccess)>
213 static const bool HasOptimizedImplementation =
false;
215 static EIGEN_DEVICE_FUNC
void run(
const Self&
self, Op& reducer,
const Device&,
typename Self::CoeffReturnType* output) {
216 const typename Self::Index num_coeffs = array_prod(
self.m_impl.dimensions());
217 *output = InnerMostDimReducer<Self, Op, Vectorizable>::reduce(
self, 0, num_coeffs, reducer);
222 #ifdef EIGEN_USE_THREADS 224 template <
typename Self,
typename Op,
225 bool Vectorizable = (Self::InputPacketAccess & Op::PacketAccess)>
226 struct FullReducerShard {
227 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void run(
const Self&
self,
typename Self::Index firstIndex,
228 typename Self::Index numValuesToReduce, Op& reducer,
229 typename Self::CoeffReturnType* output) {
230 *output = InnerMostDimReducer<Self, Op, Vectorizable>::reduce(
231 self, firstIndex, numValuesToReduce, reducer);
236 template <
typename Self,
typename Op,
bool Vectorizable>
237 struct FullReducer<Self, Op, ThreadPoolDevice, Vectorizable> {
238 static const bool HasOptimizedImplementation = !Op::IsStateful;
239 static const int PacketSize =
240 unpacket_traits<typename Self::PacketReturnType>::size;
243 static void run(
const Self&
self, Op& reducer,
const ThreadPoolDevice& device,
244 typename Self::CoeffReturnType* output) {
245 typedef typename Self::Index Index;
246 const Index num_coeffs = array_prod(
self.m_impl.dimensions());
247 if (num_coeffs == 0) {
248 *output = reducer.finalize(reducer.initialize());
251 const TensorOpCost cost =
252 self.m_impl.costPerCoeff(Vectorizable) +
253 TensorOpCost(0, 0, internal::functor_traits<Op>::Cost, Vectorizable,
255 const int num_threads = TensorCostModel<ThreadPoolDevice>::numThreads(
256 num_coeffs, cost, device.numThreads());
257 if (num_threads == 1) {
259 InnerMostDimReducer<Self, Op, Vectorizable>::reduce(
self, 0, num_coeffs, reducer);
262 const Index blocksize =
263 std::floor<Index>(
static_cast<float>(num_coeffs) / num_threads);
264 const Index numblocks = blocksize > 0 ? num_coeffs / blocksize : 0;
265 eigen_assert(num_coeffs >= numblocks * blocksize);
267 Barrier barrier(internal::convert_index<unsigned int>(numblocks));
268 MaxSizeVector<typename Self::CoeffReturnType> shards(numblocks, reducer.initialize());
269 for (Index i = 0; i < numblocks; ++i) {
270 device.enqueue_with_barrier(&barrier, &FullReducerShard<Self, Op, Vectorizable>::run,
271 self, i * blocksize, blocksize, reducer,
274 typename Self::CoeffReturnType finalShard;
275 if (numblocks * blocksize < num_coeffs) {
276 finalShard = InnerMostDimReducer<Self, Op, Vectorizable>::reduce(
277 self, numblocks * blocksize, num_coeffs - numblocks * blocksize,
280 finalShard = reducer.initialize();
284 for (Index i = 0; i < numblocks; ++i) {
285 reducer.reduce(shards[i], &finalShard);
287 *output = reducer.finalize(finalShard);
295 template <
typename Self,
typename Op,
typename Device>
296 struct InnerReducer {
297 static const bool HasOptimizedImplementation =
false;
299 EIGEN_DEVICE_FUNC
static bool run(
const Self&, Op&,
const Device&,
typename Self::CoeffReturnType*,
typename Self::Index,
typename Self::Index) {
300 eigen_assert(
false &&
"Not implemented");
306 template <
typename Self,
typename Op,
typename Device>
307 struct OuterReducer {
308 static const bool HasOptimizedImplementation =
false;
310 EIGEN_DEVICE_FUNC
static bool run(
const Self&, Op&,
const Device&,
typename Self::CoeffReturnType*,
typename Self::Index,
typename Self::Index) {
311 eigen_assert(
false &&
"Not implemented");
317 #if defined(EIGEN_USE_GPU) && defined(__CUDACC__) 318 template <
int B,
int N,
typename S,
typename R,
typename I>
319 __global__
void FullReductionKernel(R,
const S, I,
typename S::CoeffReturnType*,
unsigned int*);
322 #ifdef EIGEN_HAS_CUDA_FP16 323 template <
typename S,
typename R,
typename I>
324 __global__
void ReductionInitFullReduxKernelHalfFloat(R,
const S, I, half2*);
325 template <
int B,
int N,
typename S,
typename R,
typename I>
326 __global__
void FullReductionKernelHalfFloat(R,
const S, I, half*, half2*);
327 template <
int NPT,
typename S,
typename R,
typename I>
328 __global__
void InnerReductionKernelHalfFloat(R,
const S, I, I, half*);
332 template <
int NPT,
typename S,
typename R,
typename I>
333 __global__
void InnerReductionKernel(R,
const S, I, I,
typename S::CoeffReturnType*);
335 template <
int NPT,
typename S,
typename R,
typename I>
336 __global__
void OuterReductionKernel(R,
const S, I, I,
typename S::CoeffReturnType*);
342 template <
typename Op,
typename Dims,
typename XprType>
343 class TensorReductionOp :
public TensorBase<TensorReductionOp<Op, Dims, XprType>, ReadOnlyAccessors> {
345 typedef typename Eigen::internal::traits<TensorReductionOp>::Scalar Scalar;
346 typedef typename Eigen::NumTraits<Scalar>::Real RealScalar;
347 typedef typename internal::remove_const<typename XprType::CoeffReturnType>::type CoeffReturnType;
348 typedef typename Eigen::internal::nested<TensorReductionOp>::type Nested;
349 typedef typename Eigen::internal::traits<TensorReductionOp>::StorageKind StorageKind;
350 typedef typename Eigen::internal::traits<TensorReductionOp>::Index Index;
352 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
353 TensorReductionOp(
const XprType& expr,
const Dims& dims) : m_expr(expr), m_dims(dims)
355 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
356 TensorReductionOp(
const XprType& expr,
const Dims& dims,
const Op& reducer) : m_expr(expr), m_dims(dims), m_reducer(reducer)
359 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
360 const XprType& expression()
const {
return m_expr; }
361 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
362 const Dims& dims()
const {
return m_dims; }
363 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
364 const Op& reducer()
const {
return m_reducer; }
367 typename XprType::Nested m_expr;
374 template<
typename Op,
typename Dims,
typename ArgType,
typename Device>
375 struct TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType>, Device>
377 typedef TensorReductionOp<Op, Dims, ArgType> XprType;
378 typedef typename XprType::Index Index;
379 typedef typename TensorEvaluator<ArgType, Device>::Dimensions InputDimensions;
380 static const int NumInputDims = internal::array_size<InputDimensions>::value;
381 static const int NumReducedDims = internal::array_size<Dims>::value;
382 static const int NumOutputDims = NumInputDims - NumReducedDims;
383 typedef typename internal::conditional<NumOutputDims==0, Sizes<>, DSizes<Index, NumOutputDims> >::type Dimensions;
384 typedef typename XprType::Scalar Scalar;
385 typedef TensorEvaluator<const TensorReductionOp<Op, Dims, ArgType>, Device> Self;
386 static const bool InputPacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess;
387 typedef typename internal::remove_const<typename XprType::CoeffReturnType>::type CoeffReturnType;
388 typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
389 static const int PacketSize = internal::unpacket_traits<PacketReturnType>::size;
393 PacketAccess = Self::InputPacketAccess && Op::PacketAccess,
394 Layout = TensorEvaluator<ArgType, Device>::Layout,
399 static const bool ReducingInnerMostDims = internal::are_inner_most_dims<Dims, NumInputDims, Layout>::value;
400 static const bool PreservingInnerMostDims = internal::preserve_inner_most_dims<Dims, NumInputDims, Layout>::value;
401 static const bool RunningFullReduction = (NumOutputDims==0);
403 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator(
const XprType& op,
const Device& device)
404 : m_impl(op.expression(), device), m_reducer(op.reducer()), m_result(NULL), m_device(device)
406 EIGEN_STATIC_ASSERT((NumInputDims >= NumReducedDims), YOU_MADE_A_PROGRAMMING_MISTAKE);
407 EIGEN_STATIC_ASSERT((!ReducingInnerMostDims | !PreservingInnerMostDims | (NumReducedDims == NumInputDims)),
408 YOU_MADE_A_PROGRAMMING_MISTAKE);
411 for (
int i = 0; i < NumInputDims; ++i) {
412 m_reduced[i] =
false;
414 for (
int i = 0; i < NumReducedDims; ++i) {
415 eigen_assert(op.dims()[i] >= 0);
416 eigen_assert(op.dims()[i] < NumInputDims);
417 m_reduced[op.dims()[i]] =
true;
420 const typename TensorEvaluator<ArgType, Device>::Dimensions& input_dims = m_impl.dimensions();
421 internal::DimInitializer<Dimensions>::run(input_dims, m_reduced, &m_dimensions, &m_reducedDims);
424 if (NumOutputDims > 0) {
425 if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
426 m_outputStrides[0] = 1;
427 for (
int i = 1; i < NumOutputDims; ++i) {
428 m_outputStrides[i] = m_outputStrides[i - 1] * m_dimensions[i - 1];
431 m_outputStrides.back() = 1;
432 for (
int i = NumOutputDims - 2; i >= 0; --i) {
433 m_outputStrides[i] = m_outputStrides[i + 1] * m_dimensions[i + 1];
439 if (NumInputDims > 0) {
440 array<Index, NumInputDims> input_strides;
441 if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
442 input_strides[0] = 1;
443 for (
int i = 1; i < NumInputDims; ++i) {
444 input_strides[i] = input_strides[i-1] * input_dims[i-1];
447 input_strides.back() = 1;
448 for (
int i = NumInputDims - 2; i >= 0; --i) {
449 input_strides[i] = input_strides[i + 1] * input_dims[i + 1];
455 for (
int i = 0; i < NumInputDims; ++i) {
457 m_reducedStrides[reduceIndex] = input_strides[i];
460 m_preservedStrides[outputIndex] = input_strides[i];
467 if (NumOutputDims == 0) {
468 m_preservedStrides[0] = internal::array_prod(input_dims);
472 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
const Dimensions& dimensions()
const {
return m_dimensions; }
474 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC
bool evalSubExprsIfNeeded(CoeffReturnType* data) {
475 m_impl.evalSubExprsIfNeeded(NULL);
478 if (RunningFullReduction &&
479 internal::FullReducer<Self, Op, Device>::HasOptimizedImplementation &&
480 ((RunningOnGPU && (m_device.majorDeviceVersion() >= 3)) ||
482 bool need_assign =
false;
484 m_result =
static_cast<CoeffReturnType*
>(m_device.allocate(
sizeof(CoeffReturnType)));
489 Op reducer(m_reducer);
490 internal::FullReducer<Self, Op, Device>::run(*
this, reducer, m_device, data);
495 else if (RunningOnGPU && (m_device.majorDeviceVersion() >= 3)) {
496 bool reducing_inner_dims =
true;
497 for (
int i = 0; i < NumReducedDims; ++i) {
498 if (static_cast<int>(Layout) ==
static_cast<int>(ColMajor)) {
499 reducing_inner_dims &= m_reduced[i];
501 reducing_inner_dims &= m_reduced[NumInputDims - 1 - i];
504 if (internal::InnerReducer<Self, Op, Device>::HasOptimizedImplementation &&
505 (reducing_inner_dims || ReducingInnerMostDims)) {
506 const Index num_values_to_reduce = internal::array_prod(m_reducedDims);
507 const Index num_coeffs_to_preserve = internal::array_prod(m_dimensions);
508 if (!data && num_coeffs_to_preserve < 1024 && num_values_to_reduce > num_coeffs_to_preserve) {
509 data =
static_cast<CoeffReturnType*
>(m_device.allocate(
sizeof(CoeffReturnType) * num_coeffs_to_preserve));
512 Op reducer(m_reducer);
513 return internal::InnerReducer<Self, Op, Device>::run(*
this, reducer, m_device, data, num_values_to_reduce, num_coeffs_to_preserve) || (m_result != NULL);
516 bool preserving_inner_dims =
true;
517 for (
int i = 0; i < NumReducedDims; ++i) {
518 if (static_cast<int>(Layout) ==
static_cast<int>(ColMajor)) {
519 preserving_inner_dims &= m_reduced[NumInputDims - 1 - i];
521 preserving_inner_dims &= m_reduced[i];
524 if (internal::OuterReducer<Self, Op, Device>::HasOptimizedImplementation &&
525 preserving_inner_dims) {
526 const Index num_values_to_reduce = internal::array_prod(m_reducedDims);
527 const Index num_coeffs_to_preserve = internal::array_prod(m_dimensions);
528 if (!data && num_coeffs_to_preserve < 1024 && num_values_to_reduce > num_coeffs_to_preserve) {
529 data =
static_cast<CoeffReturnType*
>(m_device.allocate(
sizeof(CoeffReturnType) * num_coeffs_to_preserve));
532 Op reducer(m_reducer);
533 return internal::OuterReducer<Self, Op, Device>::run(*
this, reducer, m_device, data, num_values_to_reduce, num_coeffs_to_preserve) || (m_result != NULL);
539 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void cleanup() {
542 m_device.deallocate(m_result);
546 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index)
const 548 if ((RunningFullReduction || RunningOnGPU) && m_result) {
549 return *(m_result + index);
551 Op reducer(m_reducer);
552 if (ReducingInnerMostDims || RunningFullReduction) {
553 const Index num_values_to_reduce =
554 (
static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1];
555 return internal::InnerMostDimReducer<Self, Op>::reduce(*
this, firstInput(index),
556 num_values_to_reduce, reducer);
558 typename Self::CoeffReturnType accum = reducer.initialize();
559 internal::GenericDimReducer<NumReducedDims-1, Self, Op>::reduce(*
this, firstInput(index), reducer, &accum);
560 return reducer.finalize(accum);
565 template<
int LoadMode>
566 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType packet(Index index)
const 568 EIGEN_STATIC_ASSERT((PacketSize > 1), YOU_MADE_A_PROGRAMMING_MISTAKE)
569 eigen_assert(index + PacketSize - 1 < Index(internal::array_prod(dimensions())));
571 if (RunningOnGPU && m_result) {
572 return internal::pload<PacketReturnType>(m_result + index);
575 EIGEN_ALIGN_MAX
typename internal::remove_const<CoeffReturnType>::type values[PacketSize];
576 if (ReducingInnerMostDims) {
577 const Index num_values_to_reduce =
578 (
static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_preservedStrides[0] : m_preservedStrides[NumPreservedStrides - 1];
579 const Index firstIndex = firstInput(index);
580 for (Index i = 0; i < PacketSize; ++i) {
581 Op reducer(m_reducer);
582 values[i] = internal::InnerMostDimReducer<Self, Op>::reduce(*
this, firstIndex + i * num_values_to_reduce,
583 num_values_to_reduce, reducer);
585 }
else if (PreservingInnerMostDims) {
586 const Index firstIndex = firstInput(index);
587 const int innermost_dim = (
static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? 0 : NumOutputDims - 1;
589 if (((firstIndex % m_dimensions[innermost_dim]) + PacketSize - 1) < m_dimensions[innermost_dim]) {
590 Op reducer(m_reducer);
591 typename Self::PacketReturnType accum = reducer.template initializePacket<typename Self::PacketReturnType>();
592 internal::InnerMostDimPreserver<NumReducedDims-1, Self, Op>::reduce(*
this, firstIndex, reducer, &accum);
593 return reducer.finalizePacket(accum);
595 for (
int i = 0; i < PacketSize; ++i) {
596 values[i] = coeff(index + i);
600 for (
int i = 0; i < PacketSize; ++i) {
601 values[i] = coeff(index + i);
604 PacketReturnType rslt = internal::pload<PacketReturnType>(values);
609 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(
bool vectorized)
const {
610 if (RunningFullReduction && m_result) {
611 return TensorOpCost(
sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize);
613 const Index num_values_to_reduce = internal::array_prod(m_reducedDims);
614 const double compute_cost = num_values_to_reduce * internal::functor_traits<Op>::Cost;
615 return m_impl.costPerCoeff(vectorized) * num_values_to_reduce +
616 TensorOpCost(0, 0, compute_cost, vectorized, PacketSize);
620 EIGEN_DEVICE_FUNC Scalar* data()
const {
return NULL; }
623 template <
int,
typename,
typename>
friend struct internal::GenericDimReducer;
624 template <
typename,
typename,
bool>
friend struct internal::InnerMostDimReducer;
625 template <
int,
typename,
typename,
bool>
friend struct internal::InnerMostDimPreserver;
626 template <
typename S,
typename O,
typename D,
bool V>
friend struct internal::FullReducer;
627 #ifdef EIGEN_USE_THREADS 628 template <
typename S,
typename O,
bool V>
friend struct internal::FullReducerShard;
630 #if defined(EIGEN_USE_GPU) && defined(__CUDACC__) 631 template <
int B,
int N,
typename S,
typename R,
typename I>
friend void internal::FullReductionKernel(R,
const S, I,
typename S::CoeffReturnType*,
unsigned int*);
632 #ifdef EIGEN_HAS_CUDA_FP16 633 template <
typename S,
typename R,
typename I>
friend void internal::ReductionInitFullReduxKernelHalfFloat(R,
const S, I, half2*);
634 template <
int B,
int N,
typename S,
typename R,
typename I>
friend void internal::FullReductionKernelHalfFloat(R,
const S, I, half*, half2*);
635 template <
int NPT,
typename S,
typename R,
typename I>
friend void internal::InnerReductionKernelHalfFloat(R,
const S, I, I, half*);
637 template <
int NPT,
typename S,
typename R,
typename I>
friend void internal::InnerReductionKernel(R,
const S, I, I,
typename S::CoeffReturnType*);
639 template <
int NPT,
typename S,
typename R,
typename I>
friend void internal::OuterReductionKernel(R,
const S, I, I,
typename S::CoeffReturnType*);
642 template <
typename S,
typename O,
typename D>
friend struct internal::InnerReducer;
646 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Index firstInput(Index index)
const {
647 if (ReducingInnerMostDims) {
648 if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
649 return index * m_preservedStrides[0];
651 return index * m_preservedStrides[NumPreservedStrides - 1];
655 Index startInput = 0;
656 if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) {
657 for (
int i = NumOutputDims - 1; i > 0; --i) {
659 const Index idx = index / m_outputStrides[i];
660 startInput += idx * m_preservedStrides[i];
661 index -= idx * m_outputStrides[i];
663 if (PreservingInnerMostDims) {
664 eigen_assert(m_preservedStrides[0] == 1);
667 startInput += index * m_preservedStrides[0];
670 for (
int i = 0; i < NumOutputDims - 1; ++i) {
672 const Index idx = index / m_outputStrides[i];
673 startInput += idx * m_preservedStrides[i];
674 index -= idx * m_outputStrides[i];
676 if (PreservingInnerMostDims) {
677 eigen_assert(m_preservedStrides[NumPreservedStrides - 1] == 1);
680 startInput += index * m_preservedStrides[NumPreservedStrides - 1];
687 array<bool, NumInputDims> m_reduced;
689 Dimensions m_dimensions;
691 array<Index, NumOutputDims> m_outputStrides;
694 static const int NumPreservedStrides = max_n_1<NumOutputDims>::size;
695 array<Index, NumPreservedStrides> m_preservedStrides;
699 array<Index, NumReducedDims> m_reducedStrides;
702 array<Index, NumReducedDims> m_reducedDims;
705 TensorEvaluator<ArgType, Device> m_impl;
711 #if defined(EIGEN_USE_GPU) && defined(__CUDACC__) 712 static const bool RunningOnGPU = internal::is_same<Device, Eigen::GpuDevice>::value;
714 static const bool RunningOnGPU =
false;
716 CoeffReturnType* m_result;
718 const Device& m_device;
723 #endif // EIGEN_CXX11_TENSOR_TENSOR_REDUCTION_H Namespace containing all symbols from the Eigen library.
Definition: AdolcForward:45