10#ifndef EIGEN_CXX11_TENSOR_TENSOR_CONTRACTION_MAPPER_H
11#define EIGEN_CXX11_TENSOR_TENSOR_CONTRACTION_MAPPER_H
27template <
typename Tensor,
bool HasRawAccess,
template <
class>
class MakePointer_ = MakePointer>
30template <
typename Scalar,
typename Index,
int side,
typename Tensor,
31 typename nocontract_t,
typename contract_t,
int packet_size,
32 bool inner_dim_contiguous,
bool inner_dim_reordered,
int Alignment,
33 template <
class>
class MakePointer_ = MakePointer>
34class BaseTensorContractionMapper;
36template <
typename Tensor,
bool HasRawAccess,
template <
class>
class MakePointer_>
42 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE CoeffLoader(
const Tensor& tensor) : m_tensor(tensor) { }
44 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
void offsetBuffer(
typename Tensor::Index) {
45 eigen_assert(
false &&
"unsupported");
48 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
const typename MakePointer_<const typename Tensor::Scalar>::Type
50 eigen_assert(
false &&
"unsupported");
54 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
typename Tensor::Scalar coeff(
typename Tensor::Index index)
const {
return m_tensor.coeff(index); }
56 template<
int LoadMode> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
57 typename Tensor::PacketReturnType packet(
typename Tensor::Index index)
const
59 return m_tensor.template packet<LoadMode>(index);
64 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void bind(cl::sycl::handler &cgh)
const {
70 const Tensor m_tensor;
73template <
typename Tensor,
template <
class>
class MakePointer_>
74struct CoeffLoader<Tensor, true, MakePointer_> {
79 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE CoeffLoader(
const Tensor& tensor) : m_data(tensor.data()) {}
81 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
void offsetBuffer(
typename Tensor::Index offset) {
85 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
const typename MakePointer_<const typename Tensor::Scalar>::Type
90 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
typename Tensor::Scalar coeff(
typename Tensor::Index index)
const {
return loadConstant(m_data+index); }
92 template<
int LoadMode> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
93 typename Tensor::PacketReturnType packet(
typename Tensor::Index index)
const
95 return internal::ploadt_ro<typename Tensor::PacketReturnType, LoadMode>(m_data + index);
100 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void bind(cl::sycl::handler &cgh)
const {
105 typedef typename Tensor::Scalar Scalar;
107 typename MakePointer_<const Scalar>::Type m_data;
110template<
typename Scalar,
typename Index,
int side,
112 typename nocontract_t,
typename contract_t,
113 int packet_size,
bool inner_dim_contiguous,
int Alignment,
template <
class>
class MakePointer_ = MakePointer>
114class SimpleTensorContractionMapper {
117 SimpleTensorContractionMapper(
const Tensor& tensor,
118 const nocontract_t& nocontract_strides,
119 const nocontract_t& ij_strides,
120 const contract_t& contract_strides,
121 const contract_t& k_strides) :
123 m_nocontract_strides(nocontract_strides),
124 m_ij_strides(ij_strides),
125 m_contract_strides(contract_strides),
126 m_k_strides(k_strides) { }
129 DirectOffsets = CoeffLoader<Tensor, Tensor::RawAccess, MakePointer_>::DirectOffsets
132 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
void offsetBuffer(
typename Tensor::Index offset) {
133 m_tensor.offsetBuffer(offset);
137 EIGEN_STRONG_INLINE
void prefetch(
Index ) { }
140 EIGEN_STRONG_INLINE Scalar operator()(
Index row)
const {
142 return operator()(row, 0);
146 EIGEN_STRONG_INLINE Scalar operator()(
Index row,
Index col)
const {
147 return m_tensor.coeff(computeIndex(row, col));
152 const bool left = (side == Lhs);
153 EIGEN_UNUSED_VARIABLE(left);
154 Index nocontract_val = left ? row : col;
157 for (
int i =
static_cast<int>(array_size<nocontract_t>::value) - 1; i > 0; i--) {
158 const Index idx = nocontract_val / m_ij_strides[i];
159 linidx += idx * m_nocontract_strides[i];
160 nocontract_val -= idx * m_ij_strides[i];
162 if (array_size<typename Tensor::Dimensions>::value > array_size<contract_t>::value) {
163 if (side == Lhs && inner_dim_contiguous) {
164 eigen_assert(m_nocontract_strides[0] == 1);
165 linidx += nocontract_val;
167 linidx += nocontract_val * m_nocontract_strides[0];
171 Index contract_val = left ? col : row;
172 if(array_size<contract_t>::value > 0) {
174 for (
int i =
static_cast<int>(array_size<contract_t>::value) - 1; i > 0; i--) {
175 const Index idx = contract_val / m_k_strides[i];
176 linidx += idx * m_contract_strides[i];
177 contract_val -= idx * m_k_strides[i];
180 if (side == Rhs && inner_dim_contiguous) {
181 eigen_assert(m_contract_strides[0] == 1);
182 linidx += contract_val;
184 linidx += contract_val * m_contract_strides[0];
192 EIGEN_STRONG_INLINE IndexPair<Index> computeIndexPair(
Index row,
Index col,
const Index distance)
const {
193 const bool left = (side == Lhs);
194 EIGEN_UNUSED_VARIABLE(left);
195 Index nocontract_val[2] = {left ? row : col, left ? row + distance : col};
196 Index linidx[2] = {0, 0};
197 if (array_size<typename Tensor::Dimensions>::value > array_size<contract_t>::value) {
199 for (
int i =
static_cast<int>(array_size<nocontract_t>::value) - 1; i > 0; i--) {
200 const Index idx0 = nocontract_val[0] / m_ij_strides[i];
201 const Index idx1 = nocontract_val[1] / m_ij_strides[i];
202 linidx[0] += idx0 * m_nocontract_strides[i];
203 linidx[1] += idx1 * m_nocontract_strides[i];
204 nocontract_val[0] -= idx0 * m_ij_strides[i];
205 nocontract_val[1] -= idx1 * m_ij_strides[i];
207 if (side == Lhs && inner_dim_contiguous) {
208 eigen_assert(m_nocontract_strides[0] == 1);
209 linidx[0] += nocontract_val[0];
210 linidx[1] += nocontract_val[1];
212 linidx[0] += nocontract_val[0] * m_nocontract_strides[0];
213 linidx[1] += nocontract_val[1] * m_nocontract_strides[0];
217 Index contract_val[2] = {left ? col : row, left ? col : row + distance};
218 if (array_size<contract_t>::value> 0) {
220 for (
int i =
static_cast<int>(array_size<contract_t>::value) - 1; i > 0; i--) {
221 const Index idx0 = contract_val[0] / m_k_strides[i];
222 const Index idx1 = contract_val[1] / m_k_strides[i];
223 linidx[0] += idx0 * m_contract_strides[i];
224 linidx[1] += idx1 * m_contract_strides[i];
225 contract_val[0] -= idx0 * m_k_strides[i];
226 contract_val[1] -= idx1 * m_k_strides[i];
229 if (side == Rhs && inner_dim_contiguous) {
230 eigen_assert(m_contract_strides[0] == 1);
231 linidx[0] += contract_val[0];
232 linidx[1] += contract_val[1];
234 linidx[0] += contract_val[0] * m_contract_strides[0];
235 linidx[1] += contract_val[1] * m_contract_strides[0];
238 return IndexPair<Index>(linidx[0], linidx[1]);
241 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
Index firstAligned(
Index size)
const {
245 return (Alignment ==
Aligned) && (side == Lhs) && inner_dim_contiguous ? 0 : size;
247 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
Index stride()
const {
248 return ((side == Lhs) && inner_dim_contiguous && array_size<contract_t>::value > 0) ? m_contract_strides[0] : 1;
251 #ifdef EIGEN_USE_SYCL
253 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void bind(cl::sycl::handler &cgh)
const {
258 const CoeffLoader<Tensor, Tensor::RawAccess, MakePointer_>& tensor()
const {
262 const nocontract_t& nocontract_strides()
const {
263 return m_nocontract_strides;
265 const nocontract_t& ij_strides()
const {
return m_ij_strides; }
266 const contract_t& contract_strides()
const {
return m_contract_strides; }
267 const contract_t& k_strides()
const {
return m_k_strides; }
270 CoeffLoader<Tensor, Tensor::RawAccess, MakePointer_> m_tensor;
271 const nocontract_t m_nocontract_strides;
272 const nocontract_t m_ij_strides;
273 const contract_t m_contract_strides;
274 const contract_t m_k_strides;
277template<
typename Scalar,
typename Index,
int side,
279 typename nocontract_t,
typename contract_t,
280 int packet_size,
bool inner_dim_contiguous,
281 bool inner_dim_reordered,
int Alignment,
template <
class>
class MakePointer_>
282class BaseTensorContractionMapper :
public SimpleTensorContractionMapper<Scalar, Index, side, Tensor, nocontract_t, contract_t, packet_size, inner_dim_contiguous, Alignment, MakePointer_>
285 typedef SimpleTensorContractionMapper<Scalar, Index, side, Tensor, nocontract_t, contract_t, packet_size, inner_dim_contiguous, Alignment, MakePointer_> ParentMapper;
288 BaseTensorContractionMapper(
const Tensor& tensor,
289 const nocontract_t& nocontract_strides,
290 const nocontract_t& ij_strides,
291 const contract_t& contract_strides,
292 const contract_t& k_strides) :
293 ParentMapper(tensor, nocontract_strides, ij_strides, contract_strides, k_strides) { }
295 template <
typename PacketT,
int AlignmentType>
296 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
297 typename internal::enable_if<internal::unpacket_traits<PacketT>::size==packet_size,PacketT>::type
304 EIGEN_STATIC_ASSERT(packet_size % 2 == 0, YOU_MADE_A_PROGRAMMING_MISTAKE);
306 if (Tensor::PacketAccess && inner_dim_contiguous && !inner_dim_reordered) {
307 const Index index = this->computeIndex(i, j);
308 eigen_assert(this->computeIndex(i+packet_size-1, j) == index + packet_size-1);
309 return this->m_tensor.template packet<AlignmentType>(index);
312 const IndexPair<Index> indexPair = this->computeIndexPair(i, j, packet_size - 1);
313 const Index first = indexPair.first;
314 const Index lastIdx = indexPair.second;
320 if (Tensor::PacketAccess &&
321 (side == Lhs || internal::array_size<contract_t>::value <= 1 || !inner_dim_reordered) &&
322 (lastIdx - first) == (packet_size - 1)) {
324 return this->m_tensor.template packet<AlignmentType>(first);
327 EIGEN_ALIGN_MAX Scalar data[packet_size];
329 data[0] = this->m_tensor.coeff(first);
331 for (
Index k = 1; k < packet_size - 1; k += 2) {
332 const IndexPair<Index> internal_pair = this->computeIndexPair(i + k, j, 1);
333 data[k] = this->m_tensor.coeff(internal_pair.first);
334 data[k + 1] = this->m_tensor.coeff(internal_pair.second);
336 data[packet_size - 1] = this->m_tensor.coeff(lastIdx);
338 return pload<PacketT>(data);
341 template <
typename PacketT,
int AlignmentType>
342 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
343 typename internal::enable_if<internal::unpacket_traits<PacketT>::size!=packet_size,PacketT>::type
346 const Index requested_packet_size = internal::unpacket_traits<PacketT>::size;
347 EIGEN_ALIGN_MAX Scalar data[requested_packet_size];
349 const IndexPair<Index> indexPair = this->computeIndexPair(i, j, requested_packet_size - 1);
350 const Index first = indexPair.first;
351 const Index lastIdx = indexPair.second;
353 data[0] = this->m_tensor.coeff(first);
354 for (
Index k = 1; k < requested_packet_size - 1; k += 2) {
355 const IndexPair<Index> internal_pair = this->computeIndexPair(i + k, j, 1);
356 data[k] = this->m_tensor.coeff(internal_pair.first);
357 data[k + 1] = this->m_tensor.coeff(internal_pair.second);
359 data[requested_packet_size - 1] = this->m_tensor.coeff(lastIdx);
361 return pload<PacketT>(data);
364 template <
typename PacketT,
int AlignmentType>
366 EIGEN_STRONG_INLINE PacketT loadPacket(
Index i,
Index j)
const {
367 return this->load<PacketT,AlignmentType>(i,j);
372template<
typename Scalar,
typename Index,
int side,
374 typename nocontract_t,
typename contract_t,
375 bool inner_dim_contiguous,
376 bool inner_dim_reordered,
int Alignment,
template <
class>
class MakePointer_>
377class BaseTensorContractionMapper<Scalar,
Index, side, Tensor, nocontract_t, contract_t, 1, inner_dim_contiguous, inner_dim_reordered, Alignment, MakePointer_>
378 :
public SimpleTensorContractionMapper<Scalar, Index, side, Tensor, nocontract_t, contract_t, 1, inner_dim_contiguous, Alignment, MakePointer_>
381 typedef SimpleTensorContractionMapper<Scalar, Index, side, Tensor, nocontract_t, contract_t, 1, inner_dim_contiguous, Alignment, MakePointer_> ParentMapper;
384 BaseTensorContractionMapper(
const Tensor& tensor,
385 const nocontract_t& nocontract_strides,
386 const nocontract_t& ij_strides,
387 const contract_t& contract_strides,
388 const contract_t& k_strides) :
389 ParentMapper(tensor, nocontract_strides, ij_strides, contract_strides, k_strides) { }
391 template <
typename PacketT,
int> EIGEN_DEVICE_FUNC
392 EIGEN_STRONG_INLINE PacketT loadPacket(
Index i,
Index j)
const {
393 EIGEN_ALIGN_MAX Scalar data[1];
394 data[0] = this->m_tensor.coeff(this->computeIndex(i, j));
395 return pload<PacketT>(data);
397 template <
typename PacketT,
int> EIGEN_DEVICE_FUNC
398 EIGEN_STRONG_INLINE PacketT load(
Index i,
Index j)
const {
399 EIGEN_ALIGN_MAX Scalar data[1];
400 data[0] = this->m_tensor.coeff(this->computeIndex(i, j));
401 return pload<PacketT>(data);
406template<
typename Scalar,
typename Index,
int side,
408 typename nocontract_t,
typename contract_t,
410 bool inner_dim_contiguous,
bool inner_dim_reordered,
int Alignment,
template <
class>
class MakePointer_=MakePointer>
411class TensorContractionSubMapper {
414 typedef BaseTensorContractionMapper<Scalar, Index, side, Tensor, nocontract_t, contract_t, packet_size, inner_dim_contiguous, inner_dim_reordered, Alignment, MakePointer_> ParentMapper;
415 typedef TensorContractionSubMapper<Scalar, Index, side, Tensor, nocontract_t, contract_t, packet_size, inner_dim_contiguous, inner_dim_reordered, Alignment, MakePointer_> Self;
416 typedef Self LinearMapper;
421 UseDirectOffsets = ParentMapper::DirectOffsets && (side == Lhs) && inner_dim_contiguous && (array_size<contract_t>::value > 0)
424 EIGEN_DEVICE_FUNC TensorContractionSubMapper(
const ParentMapper& base_mapper,
Index vert_offset,
Index horiz_offset)
425 : m_base_mapper(base_mapper), m_vert_offset(vert_offset), m_horiz_offset(horiz_offset) {
428 if (UseDirectOffsets) {
429 Index stride = m_base_mapper.stride();
430 m_base_mapper.offsetBuffer(vert_offset + horiz_offset * stride);
434 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Scalar operator()(
Index i)
const {
435 if (UseDirectOffsets) {
436 return m_base_mapper(i, 0);
438 return m_base_mapper(i + m_vert_offset, m_horiz_offset);
440 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Scalar operator()(
Index i,
Index j)
const {
441 if (UseDirectOffsets) {
442 return m_base_mapper(i, j);
444 return m_base_mapper(i + m_vert_offset, j + m_horiz_offset);
447 template <
typename PacketT>
448 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE PacketT loadPacket(
Index i)
const {
449 if (UseDirectOffsets) {
450 return m_base_mapper.template loadPacket<PacketT,Alignment>(i, 0);
452 return m_base_mapper.template loadPacket<PacketT,Alignment>(i + m_vert_offset, m_horiz_offset);
455 template <
typename PacketT>
456 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE PacketT loadPacket(
Index i,
Index j)
const {
457 if (UseDirectOffsets) {
458 return m_base_mapper.template loadPacket<PacketT,Alignment>(i, j);
460 return m_base_mapper.template loadPacket<PacketT,Alignment>(i + m_vert_offset, j + m_horiz_offset);
463 template <
typename PacketT,
int AlignmentType>
464 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE PacketT loadPacket(
Index i,
Index j)
const {
465 if (UseDirectOffsets) {
466 return m_base_mapper.template load<PacketT,AlignmentType>(i, j);
468 return m_base_mapper.template loadPacket<PacketT,AlignmentType>(i + m_vert_offset, j + m_horiz_offset);
471 template <
typename PacketT>
472 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
void storePacket(
Index i,
const PacketT& p)
const {
473 if (UseDirectOffsets) {
474 m_base_mapper.storePacket(i, 0, p);
476 m_base_mapper.storePacket(i + m_vert_offset, m_horiz_offset, p);
479 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE LinearMapper getLinearMapper(
Index i,
Index j)
const {
480 if (UseDirectOffsets) {
481 return LinearMapper(m_base_mapper, i, j);
483 return LinearMapper(m_base_mapper, i + m_vert_offset, j + m_horiz_offset);
486 template <
typename PacketT,
int AlignmentType>
487 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE PacketT load(
Index i)
const {
488 EIGEN_STATIC_ASSERT((internal::is_same<PacketT, PacketT>::value), YOU_MADE_A_PROGRAMMING_MISTAKE);
490 if (UseDirectOffsets) {
491 return m_base_mapper.template loadPacket<PacketT,ActualAlignment>(i, 0);
493 return m_base_mapper.template loadPacket<PacketT,ActualAlignment>(i + m_vert_offset, m_horiz_offset);
496 template <
typename PacketT>
497 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
bool aligned(
Index)
const {
501 #ifdef EIGEN_USE_SYCL
503 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void bind(cl::sycl::handler &cgh)
const {
504 m_base_mapper.bind(cgh);
508 const ParentMapper& base_mapper()
const {
return m_base_mapper; }
509 Index vert_offset()
const {
return m_vert_offset; }
510 Index horiz_offset()
const {
return m_horiz_offset; }
513 ParentMapper m_base_mapper;
514 const Index m_vert_offset;
515 const Index m_horiz_offset;
519template<
typename Scalar_,
typename Index,
int side,
521 typename nocontract_t,
typename contract_t,
523 bool inner_dim_contiguous,
bool inner_dim_reordered,
int Alignment,
template <
class>
class MakePointer_=MakePointer>
524class TensorContractionInputMapper
525 :
public BaseTensorContractionMapper<Scalar_, Index, side, Tensor, nocontract_t, contract_t, packet_size, inner_dim_contiguous, inner_dim_reordered, Alignment, MakePointer_> {
528 typedef Scalar_ Scalar;
529 typedef BaseTensorContractionMapper<Scalar, Index, side, Tensor, nocontract_t, contract_t, packet_size, inner_dim_contiguous, inner_dim_reordered, Alignment, MakePointer_> Base;
530 typedef TensorContractionSubMapper<Scalar, Index, side, Tensor, nocontract_t, contract_t, packet_size, inner_dim_contiguous, inner_dim_reordered, Alignment, MakePointer_> SubMapper;
531 typedef SubMapper VectorMapper;
533 EIGEN_DEVICE_FUNC TensorContractionInputMapper(
const Tensor& tensor,
534 const nocontract_t& nocontract_strides,
535 const nocontract_t& ij_strides,
536 const contract_t& contract_strides,
537 const contract_t& k_strides)
538 : Base(tensor, nocontract_strides, ij_strides, contract_strides, k_strides) { }
541 EIGEN_STRONG_INLINE SubMapper getSubMapper(
Index i,
Index j)
const {
542 return SubMapper(*
this, i, j);
545 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE VectorMapper getVectorMapper(
Index i,
Index j)
const {
546 return VectorMapper(*
this, i, j);
549 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
const CoeffLoader<Tensor, Tensor::RawAccess, MakePointer_>& get_tensor()
const {
550 return Base::m_tensor;
555template <
typename T>
struct TensorContractionInputMapperTrait;
557template<
typename Scalar_,
typename Index_,
int side_,
559 typename nocontract_t_,
typename contract_t_,
561 bool inner_dim_contiguous_,
bool inner_dim_reordered_,
int Alignment_,
template <
class>
class MakePointer_>
562struct TensorContractionInputMapperTrait<TensorContractionInputMapper<Scalar_, Index_, side_, Tensor_,
563 nocontract_t_, contract_t_, packet_size_, inner_dim_contiguous_,
564 inner_dim_reordered_, Alignment_, MakePointer_> > {
566 typedef Tensor_ XprType;
567 static const bool inner_dim_contiguous = inner_dim_contiguous_;
568 static const bool inner_dim_reordered = inner_dim_reordered_;
Namespace containing all symbols from the Eigen library.
EIGEN_DEFAULT_DENSE_INDEX_TYPE Index