12b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// This file is part of Eigen, a lightweight C++ template library 22b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// for linear algebra. 32b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// 42b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// Copyright (C) 2014 Benoit Steiner <benoit.steiner.goog@gmail.com> 52b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// 62b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// This Source Code Form is subject to the terms of the Mozilla 72b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// Public License v. 2.0. If a copy of the MPL was not distributed 82b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// with this file, You can obtain one at http://mozilla.org/MPL/2.0/. 92b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 102b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#ifndef EIGEN_CXX11_TENSOR_TENSOR_EVALUATOR_H 112b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#define EIGEN_CXX11_TENSOR_TENSOR_EVALUATOR_H 122b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 132b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangnamespace Eigen { 142b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 152b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang/** \class TensorEvaluator 162b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang * \ingroup CXX11_Tensor_Module 172b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang * 182b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang * \brief The tensor evaluator classes. 192b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang * 202b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang * These classes are responsible for the evaluation of the tensor expression. 212b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang * 222b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang * TODO: add support for more types of expressions, in particular expressions 232b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang * leading to lvalues (slicing, reshaping, etc...) 242b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang */ 252b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 262b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// Generic evaluator 272b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<typename Derived, typename Device> 282b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangstruct TensorEvaluator 292b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang{ 302b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename Derived::Index Index; 312b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename Derived::Scalar Scalar; 322b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename Derived::Scalar CoeffReturnType; 332b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType; 342b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename Derived::Dimensions Dimensions; 352b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 362b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang // NumDimensions is -1 for variable dim tensors 372b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang static const int NumCoords = internal::traits<Derived>::NumDimensions > 0 ? 382b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang internal::traits<Derived>::NumDimensions : 0; 392b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 402b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang enum { 412b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang IsAligned = Derived::IsAligned, 422b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang PacketAccess = (internal::unpacket_traits<PacketReturnType>::size > 1), 432b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang Layout = Derived::Layout, 442b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang CoordAccess = NumCoords > 0, 452b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang RawAccess = true 462b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang }; 472b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 482b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator(const Derived& m, const Device& device) 492b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang : m_data(const_cast<typename internal::traits<Derived>::template MakePointer<Scalar>::Type>(m.data())), m_dims(m.dimensions()), m_device(device), m_impl(m) 502b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { } 512b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 522b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang // Used for accessor extraction in SYCL Managed TensorMap: 532b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const Derived& derived() const { return m_impl; } 542b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Dimensions& dimensions() const { return m_dims; } 552b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 562b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(CoeffReturnType* dest) { 572b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang if (dest) { 582b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_device.memcpy((void*)dest, m_data, sizeof(Scalar) * m_dims.TotalSize()); 592b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return false; 602b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 612b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return true; 622b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 632b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 642b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void cleanup() { } 652b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 662b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index) const { 672b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang eigen_assert(m_data); 682b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_data[index]; 692b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 702b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 712b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar& coeffRef(Index index) { 722b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang eigen_assert(m_data); 732b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_data[index]; 742b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 752b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 762b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang template<int LoadMode> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE 772b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang PacketReturnType packet(Index index) const 782b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 792b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return internal::ploadt<PacketReturnType, LoadMode>(m_data + index); 802b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 812b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 822b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang template <int StoreMode> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE 832b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang void writePacket(Index index, const PacketReturnType& x) 842b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 852b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return internal::pstoret<Scalar, PacketReturnType, StoreMode>(m_data + index, x); 862b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 872b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 882b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(const array<DenseIndex, NumCoords>& coords) const { 892b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang eigen_assert(m_data); 902b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) { 912b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_data[m_dims.IndexOfColMajor(coords)]; 922b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } else { 932b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_data[m_dims.IndexOfRowMajor(coords)]; 942b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 952b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 962b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 972b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar& coeffRef(const array<DenseIndex, NumCoords>& coords) { 982b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang eigen_assert(m_data); 992b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang if (static_cast<int>(Layout) == static_cast<int>(ColMajor)) { 1002b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_data[m_dims.IndexOfColMajor(coords)]; 1012b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } else { 1022b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_data[m_dims.IndexOfRowMajor(coords)]; 1032b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 1042b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 1052b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1062b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(bool vectorized) const { 1072b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return TensorOpCost(sizeof(CoeffReturnType), 0, 0, vectorized, 1082b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang internal::unpacket_traits<PacketReturnType>::size); 1092b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 1102b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1112b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC typename internal::traits<Derived>::template MakePointer<Scalar>::Type data() const { return m_data; } 1122b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1132b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to construct sycl buffer from raw pointer 1142b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const Device& device() const{return m_device;} 1152b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1162b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang protected: 1172b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typename internal::traits<Derived>::template MakePointer<Scalar>::Type m_data; 1182b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang Dimensions m_dims; 1192b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const Device& m_device; 1202b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const Derived& m_impl; 1212b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang}; 1222b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1232b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangnamespace { 1242b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate <typename T> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE 1252b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao WangT loadConstant(const T* address) { 1262b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return *address; 1272b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1282b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// Use the texture cache on CUDA devices whenever possible 1292b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350 1302b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE 1312b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangfloat loadConstant(const float* address) { 1322b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return __ldg(address); 1332b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1342b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE 1352b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangdouble loadConstant(const double* address) { 1362b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return __ldg(address); 1372b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1382b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE 1392b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao WangEigen::half loadConstant(const Eigen::half* address) { 1402b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return Eigen::half(half_impl::raw_uint16_to_half(__ldg(&address->x))); 1412b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1422b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#endif 1432b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1442b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1452b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1462b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// Default evaluator for rvalues 1472b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<typename Derived, typename Device> 1482b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangstruct TensorEvaluator<const Derived, Device> 1492b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang{ 1502b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename Derived::Index Index; 1512b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename Derived::Scalar Scalar; 1522b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename Derived::Scalar CoeffReturnType; 1532b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType; 1542b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename Derived::Dimensions Dimensions; 1552b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1562b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang // NumDimensions is -1 for variable dim tensors 1572b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang static const int NumCoords = internal::traits<Derived>::NumDimensions > 0 ? 1582b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang internal::traits<Derived>::NumDimensions : 0; 1592b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1602b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang enum { 1612b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang IsAligned = Derived::IsAligned, 1622b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang PacketAccess = (internal::unpacket_traits<PacketReturnType>::size > 1), 1632b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang Layout = Derived::Layout, 1642b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang CoordAccess = NumCoords > 0, 1652b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang RawAccess = true 1662b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang }; 1672b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1682b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang // Used for accessor extraction in SYCL Managed TensorMap: 1692b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const Derived& derived() const { return m_impl; } 1702b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1712b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorEvaluator(const Derived& m, const Device& device) 1722b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang : m_data(m.data()), m_dims(m.dimensions()), m_device(device), m_impl(m) 1732b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { } 1742b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1752b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Dimensions& dimensions() const { return m_dims; } 1762b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1772b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(CoeffReturnType* data) { 1782b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang if (!NumTraits<typename internal::remove_const<Scalar>::type>::RequireInitialization && data) { 1792b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_device.memcpy((void*)data, m_data, m_dims.TotalSize() * sizeof(Scalar)); 1802b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return false; 1812b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 1822b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return true; 1832b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 1842b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1852b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void cleanup() { } 1862b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1872b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index) const { 1882b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang eigen_assert(m_data); 1892b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return loadConstant(m_data+index); 1902b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 1912b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1922b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang template<int LoadMode> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE 1932b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang PacketReturnType packet(Index index) const 1942b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 1952b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return internal::ploadt_ro<PacketReturnType, LoadMode>(m_data + index); 1962b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 1972b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1982b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(const array<DenseIndex, NumCoords>& coords) const { 1992b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang eigen_assert(m_data); 2002b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const Index index = (static_cast<int>(Layout) == static_cast<int>(ColMajor)) ? m_dims.IndexOfColMajor(coords) 2012b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang : m_dims.IndexOfRowMajor(coords); 2022b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return loadConstant(m_data+index); 2032b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 2042b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2052b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(bool vectorized) const { 2062b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return TensorOpCost(sizeof(CoeffReturnType), 0, 0, vectorized, 2072b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang internal::unpacket_traits<PacketReturnType>::size); 2082b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 2092b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2102b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC typename internal::traits<Derived>::template MakePointer<const Scalar>::Type data() const { return m_data; } 2112b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2122b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// added for sycl in order to construct the buffer from the sycl device 2132b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const Device& device() const{return m_device;} 2142b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2152b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang protected: 2162b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typename internal::traits<Derived>::template MakePointer<const Scalar>::Type m_data; 2172b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang Dimensions m_dims; 2182b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const Device& m_device; 2192b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const Derived& m_impl; 2202b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang}; 2212b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2222b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2232b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2242b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2252b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// -------------------- CwiseNullaryOp -------------------- 2262b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2272b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<typename NullaryOp, typename ArgType, typename Device> 2282b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangstruct TensorEvaluator<const TensorCwiseNullaryOp<NullaryOp, ArgType>, Device> 2292b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang{ 2302b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef TensorCwiseNullaryOp<NullaryOp, ArgType> XprType; 2312b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2322b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang enum { 2332b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang IsAligned = true, 2342b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang PacketAccess = internal::functor_traits<NullaryOp>::PacketAccess, 2352b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang Layout = TensorEvaluator<ArgType, Device>::Layout, 2362b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang CoordAccess = false, // to be implemented 2372b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang RawAccess = false 2382b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang }; 2392b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2402b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC 2412b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorEvaluator(const XprType& op, const Device& device) 2422b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang : m_functor(op.functor()), m_argImpl(op.nestedExpression(), device), m_wrapper() 2432b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { } 2442b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2452b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename XprType::Index Index; 2462b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename XprType::Scalar Scalar; 2472b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename internal::traits<XprType>::Scalar CoeffReturnType; 2482b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType; 2492b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang static const int PacketSize = internal::unpacket_traits<PacketReturnType>::size; 2502b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename TensorEvaluator<ArgType, Device>::Dimensions Dimensions; 2512b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2522b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC const Dimensions& dimensions() const { return m_argImpl.dimensions(); } 2532b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2542b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(CoeffReturnType*) { return true; } 2552b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void cleanup() { } 2562b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2572b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC CoeffReturnType coeff(Index index) const 2582b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 2592b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_wrapper(m_functor, index); 2602b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 2612b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2622b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang template<int LoadMode> 2632b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType packet(Index index) const 2642b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 2652b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_wrapper.template packetOp<PacketReturnType, Index>(m_functor, index); 2662b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 2672b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2682b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost 2692b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang costPerCoeff(bool vectorized) const { 2702b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return TensorOpCost(sizeof(CoeffReturnType), 0, 0, vectorized, 2712b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang internal::unpacket_traits<PacketReturnType>::size); 2722b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 2732b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2742b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; } 2752b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2762b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 2772b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const TensorEvaluator<ArgType, Device>& impl() const { return m_argImpl; } 2782b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 2792b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang NullaryOp functor() const { return m_functor; } 2802b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2812b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2822b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang private: 2832b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const NullaryOp m_functor; 2842b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorEvaluator<ArgType, Device> m_argImpl; 2852b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const internal::nullary_wrapper<CoeffReturnType,NullaryOp> m_wrapper; 2862b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang}; 2872b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2882b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2892b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2902b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// -------------------- CwiseUnaryOp -------------------- 2912b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2922b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<typename UnaryOp, typename ArgType, typename Device> 2932b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangstruct TensorEvaluator<const TensorCwiseUnaryOp<UnaryOp, ArgType>, Device> 2942b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang{ 2952b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef TensorCwiseUnaryOp<UnaryOp, ArgType> XprType; 2962b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2972b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang enum { 2982b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang IsAligned = TensorEvaluator<ArgType, Device>::IsAligned, 2992b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang PacketAccess = TensorEvaluator<ArgType, Device>::PacketAccess & internal::functor_traits<UnaryOp>::PacketAccess, 3002b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang Layout = TensorEvaluator<ArgType, Device>::Layout, 3012b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang CoordAccess = false, // to be implemented 3022b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang RawAccess = false 3032b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang }; 3042b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3052b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC TensorEvaluator(const XprType& op, const Device& device) 3062b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang : m_functor(op.functor()), 3072b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_argImpl(op.nestedExpression(), device) 3082b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { } 3092b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3102b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename XprType::Index Index; 3112b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename XprType::Scalar Scalar; 3122b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename internal::traits<XprType>::Scalar CoeffReturnType; 3132b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType; 3142b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang static const int PacketSize = internal::unpacket_traits<PacketReturnType>::size; 3152b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename TensorEvaluator<ArgType, Device>::Dimensions Dimensions; 3162b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3172b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC const Dimensions& dimensions() const { return m_argImpl.dimensions(); } 3182b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3192b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(Scalar*) { 3202b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_argImpl.evalSubExprsIfNeeded(NULL); 3212b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return true; 3222b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 3232b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void cleanup() { 3242b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_argImpl.cleanup(); 3252b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 3262b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3272b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC CoeffReturnType coeff(Index index) const 3282b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 3292b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_functor(m_argImpl.coeff(index)); 3302b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 3312b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3322b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang template<int LoadMode> 3332b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType packet(Index index) const 3342b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 3352b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_functor.packetOp(m_argImpl.template packet<LoadMode>(index)); 3362b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 3372b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3382b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(bool vectorized) const { 3392b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const double functor_cost = internal::functor_traits<UnaryOp>::Cost; 3402b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_argImpl.costPerCoeff(vectorized) + 3412b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorOpCost(0, 0, functor_cost, vectorized, PacketSize); 3422b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 3432b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3442b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; } 3452b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3462b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 3472b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const TensorEvaluator<ArgType, Device> & impl() const { return m_argImpl; } 3482b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// added for sycl in order to construct the buffer from sycl device 3492b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang UnaryOp functor() const { return m_functor; } 3502b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3512b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3522b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang private: 3532b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const UnaryOp m_functor; 3542b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorEvaluator<ArgType, Device> m_argImpl; 3552b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang}; 3562b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3572b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3582b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// -------------------- CwiseBinaryOp -------------------- 3592b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3602b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<typename BinaryOp, typename LeftArgType, typename RightArgType, typename Device> 3612b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangstruct TensorEvaluator<const TensorCwiseBinaryOp<BinaryOp, LeftArgType, RightArgType>, Device> 3622b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang{ 3632b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef TensorCwiseBinaryOp<BinaryOp, LeftArgType, RightArgType> XprType; 3642b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3652b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang enum { 3662b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang IsAligned = TensorEvaluator<LeftArgType, Device>::IsAligned & TensorEvaluator<RightArgType, Device>::IsAligned, 3672b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang PacketAccess = TensorEvaluator<LeftArgType, Device>::PacketAccess & TensorEvaluator<RightArgType, Device>::PacketAccess & 3682b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang internal::functor_traits<BinaryOp>::PacketAccess, 3692b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang Layout = TensorEvaluator<LeftArgType, Device>::Layout, 3702b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang CoordAccess = false, // to be implemented 3712b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang RawAccess = false 3722b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang }; 3732b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3742b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC TensorEvaluator(const XprType& op, const Device& device) 3752b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang : m_functor(op.functor()), 3762b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_leftImpl(op.lhsExpression(), device), 3772b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_rightImpl(op.rhsExpression(), device) 3782b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 3792b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_STATIC_ASSERT((static_cast<int>(TensorEvaluator<LeftArgType, Device>::Layout) == static_cast<int>(TensorEvaluator<RightArgType, Device>::Layout) || internal::traits<XprType>::NumDimensions <= 1), YOU_MADE_A_PROGRAMMING_MISTAKE); 3802b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang eigen_assert(dimensions_match(m_leftImpl.dimensions(), m_rightImpl.dimensions())); 3812b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 3822b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3832b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename XprType::Index Index; 3842b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename XprType::Scalar Scalar; 3852b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename internal::traits<XprType>::Scalar CoeffReturnType; 3862b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType; 3872b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang static const int PacketSize = internal::unpacket_traits<PacketReturnType>::size; 3882b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename TensorEvaluator<LeftArgType, Device>::Dimensions Dimensions; 3892b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3902b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC const Dimensions& dimensions() const 3912b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 3922b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang // TODO: use right impl instead if right impl dimensions are known at compile time. 3932b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_leftImpl.dimensions(); 3942b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 3952b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3962b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(CoeffReturnType*) { 3972b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_leftImpl.evalSubExprsIfNeeded(NULL); 3982b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_rightImpl.evalSubExprsIfNeeded(NULL); 3992b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return true; 4002b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 4012b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void cleanup() { 4022b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_leftImpl.cleanup(); 4032b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_rightImpl.cleanup(); 4042b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 4052b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4062b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC CoeffReturnType coeff(Index index) const 4072b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 4082b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_functor(m_leftImpl.coeff(index), m_rightImpl.coeff(index)); 4092b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 4102b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang template<int LoadMode> 4112b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType packet(Index index) const 4122b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 4132b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_functor.packetOp(m_leftImpl.template packet<LoadMode>(index), m_rightImpl.template packet<LoadMode>(index)); 4142b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 4152b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4162b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost 4172b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang costPerCoeff(bool vectorized) const { 4182b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const double functor_cost = internal::functor_traits<BinaryOp>::Cost; 4192b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_leftImpl.costPerCoeff(vectorized) + 4202b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_rightImpl.costPerCoeff(vectorized) + 4212b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorOpCost(0, 0, functor_cost, vectorized, PacketSize); 4222b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 4232b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4242b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; } 4252b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 4262b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const TensorEvaluator<LeftArgType, Device>& left_impl() const { return m_leftImpl; } 4272b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 4282b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const TensorEvaluator<RightArgType, Device>& right_impl() const { return m_rightImpl; } 4292b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 4302b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang BinaryOp functor() const { return m_functor; } 4312b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4322b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang private: 4332b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const BinaryOp m_functor; 4342b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorEvaluator<LeftArgType, Device> m_leftImpl; 4352b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorEvaluator<RightArgType, Device> m_rightImpl; 4362b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang}; 4372b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4382b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// -------------------- CwiseTernaryOp -------------------- 4392b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4402b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<typename TernaryOp, typename Arg1Type, typename Arg2Type, typename Arg3Type, typename Device> 4412b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangstruct TensorEvaluator<const TensorCwiseTernaryOp<TernaryOp, Arg1Type, Arg2Type, Arg3Type>, Device> 4422b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang{ 4432b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef TensorCwiseTernaryOp<TernaryOp, Arg1Type, Arg2Type, Arg3Type> XprType; 4442b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4452b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang enum { 4462b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang IsAligned = TensorEvaluator<Arg1Type, Device>::IsAligned & TensorEvaluator<Arg2Type, Device>::IsAligned & TensorEvaluator<Arg3Type, Device>::IsAligned, 4472b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang PacketAccess = TensorEvaluator<Arg1Type, Device>::PacketAccess & TensorEvaluator<Arg2Type, Device>::PacketAccess & TensorEvaluator<Arg3Type, Device>::PacketAccess & 4482b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang internal::functor_traits<TernaryOp>::PacketAccess, 4492b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang Layout = TensorEvaluator<Arg1Type, Device>::Layout, 4502b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang CoordAccess = false, // to be implemented 4512b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang RawAccess = false 4522b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang }; 4532b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4542b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC TensorEvaluator(const XprType& op, const Device& device) 4552b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang : m_functor(op.functor()), 4562b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg1Impl(op.arg1Expression(), device), 4572b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg2Impl(op.arg2Expression(), device), 4582b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg3Impl(op.arg3Expression(), device) 4592b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 4602b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_STATIC_ASSERT((static_cast<int>(TensorEvaluator<Arg1Type, Device>::Layout) == static_cast<int>(TensorEvaluator<Arg3Type, Device>::Layout) || internal::traits<XprType>::NumDimensions <= 1), YOU_MADE_A_PROGRAMMING_MISTAKE); 4612b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4622b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_STATIC_ASSERT((internal::is_same<typename internal::traits<Arg1Type>::StorageKind, 4632b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typename internal::traits<Arg2Type>::StorageKind>::value), 4642b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang STORAGE_KIND_MUST_MATCH) 4652b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_STATIC_ASSERT((internal::is_same<typename internal::traits<Arg1Type>::StorageKind, 4662b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typename internal::traits<Arg3Type>::StorageKind>::value), 4672b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang STORAGE_KIND_MUST_MATCH) 4682b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_STATIC_ASSERT((internal::is_same<typename internal::traits<Arg1Type>::Index, 4692b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typename internal::traits<Arg2Type>::Index>::value), 4702b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang STORAGE_INDEX_MUST_MATCH) 4712b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_STATIC_ASSERT((internal::is_same<typename internal::traits<Arg1Type>::Index, 4722b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typename internal::traits<Arg3Type>::Index>::value), 4732b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang STORAGE_INDEX_MUST_MATCH) 4742b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4752b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang eigen_assert(dimensions_match(m_arg1Impl.dimensions(), m_arg2Impl.dimensions()) && dimensions_match(m_arg1Impl.dimensions(), m_arg3Impl.dimensions())); 4762b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 4772b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4782b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename XprType::Index Index; 4792b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename XprType::Scalar Scalar; 4802b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename internal::traits<XprType>::Scalar CoeffReturnType; 4812b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType; 4822b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang static const int PacketSize = internal::unpacket_traits<PacketReturnType>::size; 4832b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename TensorEvaluator<Arg1Type, Device>::Dimensions Dimensions; 4842b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4852b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC const Dimensions& dimensions() const 4862b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 4872b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang // TODO: use arg2 or arg3 dimensions if they are known at compile time. 4882b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_arg1Impl.dimensions(); 4892b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 4902b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 4912b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(CoeffReturnType*) { 4922b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg1Impl.evalSubExprsIfNeeded(NULL); 4932b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg2Impl.evalSubExprsIfNeeded(NULL); 4942b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg3Impl.evalSubExprsIfNeeded(NULL); 4952b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return true; 4962b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 4972b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void cleanup() { 4982b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg1Impl.cleanup(); 4992b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg2Impl.cleanup(); 5002b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg3Impl.cleanup(); 5012b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 5022b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5032b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC CoeffReturnType coeff(Index index) const 5042b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 5052b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_functor(m_arg1Impl.coeff(index), m_arg2Impl.coeff(index), m_arg3Impl.coeff(index)); 5062b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 5072b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang template<int LoadMode> 5082b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType packet(Index index) const 5092b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 5102b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_functor.packetOp(m_arg1Impl.template packet<LoadMode>(index), 5112b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg2Impl.template packet<LoadMode>(index), 5122b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg3Impl.template packet<LoadMode>(index)); 5132b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 5142b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5152b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost 5162b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang costPerCoeff(bool vectorized) const { 5172b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const double functor_cost = internal::functor_traits<TernaryOp>::Cost; 5182b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_arg1Impl.costPerCoeff(vectorized) + 5192b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg2Impl.costPerCoeff(vectorized) + 5202b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_arg3Impl.costPerCoeff(vectorized) + 5212b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorOpCost(0, 0, functor_cost, vectorized, PacketSize); 5222b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 5232b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5242b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC CoeffReturnType* data() const { return NULL; } 5252b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5262b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 5272b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const TensorEvaluator<Arg1Type, Device> & arg1Impl() const { return m_arg1Impl; } 5282b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 5292b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const TensorEvaluator<Arg2Type, Device>& arg2Impl() const { return m_arg2Impl; } 5302b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 5312b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const TensorEvaluator<Arg3Type, Device>& arg3Impl() const { return m_arg3Impl; } 5322b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5332b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang private: 5342b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const TernaryOp m_functor; 5352b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorEvaluator<Arg1Type, Device> m_arg1Impl; 5362b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorEvaluator<Arg2Type, Device> m_arg2Impl; 5372b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorEvaluator<Arg3Type, Device> m_arg3Impl; 5382b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang}; 5392b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5402b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5412b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// -------------------- SelectOp -------------------- 5422b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5432b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<typename IfArgType, typename ThenArgType, typename ElseArgType, typename Device> 5442b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangstruct TensorEvaluator<const TensorSelectOp<IfArgType, ThenArgType, ElseArgType>, Device> 5452b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang{ 5462b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef TensorSelectOp<IfArgType, ThenArgType, ElseArgType> XprType; 5472b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename XprType::Scalar Scalar; 5482b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5492b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang enum { 5502b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang IsAligned = TensorEvaluator<ThenArgType, Device>::IsAligned & TensorEvaluator<ElseArgType, Device>::IsAligned, 5512b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang PacketAccess = TensorEvaluator<ThenArgType, Device>::PacketAccess & TensorEvaluator<ElseArgType, Device>::PacketAccess & 5522b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang internal::packet_traits<Scalar>::HasBlend, 5532b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang Layout = TensorEvaluator<IfArgType, Device>::Layout, 5542b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang CoordAccess = false, // to be implemented 5552b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang RawAccess = false 5562b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang }; 5572b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5582b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC TensorEvaluator(const XprType& op, const Device& device) 5592b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang : m_condImpl(op.ifExpression(), device), 5602b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_thenImpl(op.thenExpression(), device), 5612b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_elseImpl(op.elseExpression(), device) 5622b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 5632b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_STATIC_ASSERT((static_cast<int>(TensorEvaluator<IfArgType, Device>::Layout) == static_cast<int>(TensorEvaluator<ThenArgType, Device>::Layout)), YOU_MADE_A_PROGRAMMING_MISTAKE); 5642b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_STATIC_ASSERT((static_cast<int>(TensorEvaluator<IfArgType, Device>::Layout) == static_cast<int>(TensorEvaluator<ElseArgType, Device>::Layout)), YOU_MADE_A_PROGRAMMING_MISTAKE); 5652b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang eigen_assert(dimensions_match(m_condImpl.dimensions(), m_thenImpl.dimensions())); 5662b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang eigen_assert(dimensions_match(m_thenImpl.dimensions(), m_elseImpl.dimensions())); 5672b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 5682b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5692b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename XprType::Index Index; 5702b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename internal::traits<XprType>::Scalar CoeffReturnType; 5712b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType; 5722b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang static const int PacketSize = internal::unpacket_traits<PacketReturnType>::size; 5732b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef typename TensorEvaluator<IfArgType, Device>::Dimensions Dimensions; 5742b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5752b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC const Dimensions& dimensions() const 5762b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 5772b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang // TODO: use then or else impl instead if they happen to be known at compile time. 5782b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_condImpl.dimensions(); 5792b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 5802b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5812b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(CoeffReturnType*) { 5822b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_condImpl.evalSubExprsIfNeeded(NULL); 5832b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_thenImpl.evalSubExprsIfNeeded(NULL); 5842b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_elseImpl.evalSubExprsIfNeeded(NULL); 5852b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return true; 5862b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 5872b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void cleanup() { 5882b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_condImpl.cleanup(); 5892b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_thenImpl.cleanup(); 5902b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_elseImpl.cleanup(); 5912b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 5922b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 5932b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC CoeffReturnType coeff(Index index) const 5942b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 5952b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_condImpl.coeff(index) ? m_thenImpl.coeff(index) : m_elseImpl.coeff(index); 5962b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 5972b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang template<int LoadMode> 5982b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC PacketReturnType packet(Index index) const 5992b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang { 6002b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang internal::Selector<PacketSize> select; 6012b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang for (Index i = 0; i < PacketSize; ++i) { 6022b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang select.select[i] = m_condImpl.coeff(index+i); 6032b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 6042b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return internal::pblend(select, 6052b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_thenImpl.template packet<LoadMode>(index), 6062b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_elseImpl.template packet<LoadMode>(index)); 6072b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 6082b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 6092b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost 6102b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang costPerCoeff(bool vectorized) const { 6112b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return m_condImpl.costPerCoeff(vectorized) + 6122b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang m_thenImpl.costPerCoeff(vectorized) 6132b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang .cwiseMax(m_elseImpl.costPerCoeff(vectorized)); 6142b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang } 6152b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 6162b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType* data() const { return NULL; } 6172b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 6182b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const TensorEvaluator<IfArgType, Device> & cond_impl() const { return m_condImpl; } 6192b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 6202b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const TensorEvaluator<ThenArgType, Device>& then_impl() const { return m_thenImpl; } 6212b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang /// required by sycl in order to extract the accessor 6222b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang const TensorEvaluator<ElseArgType, Device>& else_impl() const { return m_elseImpl; } 6232b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 6242b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang private: 6252b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorEvaluator<IfArgType, Device> m_condImpl; 6262b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorEvaluator<ThenArgType, Device> m_thenImpl; 6272b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang TensorEvaluator<ElseArgType, Device> m_elseImpl; 6282b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang}; 6292b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 6302b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 6312b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} // end namespace Eigen 6322b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 6332b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#endif // EIGEN_CXX11_TENSOR_TENSOR_EVALUATOR_H 634