104struct TensorEvaluator<const TensorForcedEvalOp<ArgType_>, Device>
106 typedef const typename internal::remove_all<ArgType_>::type ArgType;
107 typedef TensorForcedEvalOp<ArgType> XprType;
108 typedef typename ArgType::Scalar Scalar;
109 typedef typename TensorEvaluator<ArgType, Device>::Dimensions Dimensions;
110 typedef typename XprType::Index Index;
111 typedef typename XprType::CoeffReturnType CoeffReturnType;
112 typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
113 static const int PacketSize = PacketType<CoeffReturnType, Device>::size;
114 typedef typename Eigen::internal::traits<XprType>::PointerType TensorPointerType;
115 typedef StorageMemory<CoeffReturnType, Device> Storage;
116 typedef typename Storage::Type EvaluatorPointerType;
120 PacketAccess = (PacketType<CoeffReturnType, Device>::size > 1),
121 BlockAccess = internal::is_arithmetic<CoeffReturnType>::value,
122 PreferBlockAccess =
false,
123 Layout = TensorEvaluator<ArgType, Device>::Layout,
127 static const int NumDims = internal::traits<ArgType>::NumDimensions;
130 typedef internal::TensorBlockDescriptor<NumDims, Index> TensorBlockDesc;
131 typedef internal::TensorBlockScratchAllocator<Device> TensorBlockScratch;
133 typedef typename internal::TensorMaterializedBlock<CoeffReturnType, NumDims,
138 TensorEvaluator(
const XprType& op,
const Device& device)
139 : m_impl(op.expression(), device), m_op(op.expression()),
140 m_device(device), m_buffer(NULL)
143 EIGEN_DEVICE_FUNC
const Dimensions& dimensions()
const {
return m_impl.dimensions(); }
145 EIGEN_STRONG_INLINE
bool evalSubExprsIfNeeded(EvaluatorPointerType) {
146 const Index numValues = internal::array_prod(m_impl.dimensions());
147 m_buffer = m_device.get((CoeffReturnType*)m_device.allocate_temp(numValues *
sizeof(CoeffReturnType)));
149 internal::non_integral_type_placement_new<Device, CoeffReturnType>()(numValues, m_buffer);
151 typedef TensorEvalToOp< const typename internal::remove_const<ArgType>::type > EvalTo;
152 EvalTo evalToTmp(m_device.get(m_buffer), m_op);
154 internal::TensorExecutor<
155 const EvalTo,
typename internal::remove_const<Device>::type,
156 internal::IsVectorizable<Device, const ArgType>::value,
157 internal::IsTileable<Device, const ArgType>::value>::
158 run(evalToTmp, m_device);
163#ifdef EIGEN_USE_THREADS
164 template <
typename EvalSubExprsCallback>
165 EIGEN_STRONG_INLINE
void evalSubExprsIfNeededAsync(
166 EvaluatorPointerType, EvalSubExprsCallback done) {
167 const Index numValues = internal::array_prod(m_impl.dimensions());
168 m_buffer = m_device.get((CoeffReturnType*)m_device.allocate_temp(
169 numValues *
sizeof(CoeffReturnType)));
170 typedef TensorEvalToOp<const typename internal::remove_const<ArgType>::type>
172 EvalTo evalToTmp(m_device.get(m_buffer), m_op);
174 auto on_done = std::bind([](EvalSubExprsCallback done_) { done_(
true); },
176 internal::TensorAsyncExecutor<
177 const EvalTo,
typename internal::remove_const<Device>::type,
179 internal::IsVectorizable<Device, const ArgType>::value,
180 internal::IsTileable<Device, const ArgType>::value>::
181 runAsync(evalToTmp, m_device, std::move(on_done));
185 EIGEN_STRONG_INLINE
void cleanup() {
186 m_device.deallocate_temp(m_buffer);
190 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index)
const
192 return m_buffer[index];
195 template<
int LoadMode>
196 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE PacketReturnType packet(Index index)
const
198 return internal::ploadt<PacketReturnType, LoadMode>(m_buffer + index);
201 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
202 internal::TensorBlockResourceRequirements getResourceRequirements()
const {
203 return internal::TensorBlockResourceRequirements::any();
206 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorBlock
207 block(TensorBlockDesc& desc, TensorBlockScratch& scratch,
208 bool =
false)
const {
209 assert(m_buffer != NULL);
210 return TensorBlock::materialize(m_buffer, m_impl.dimensions(), desc, scratch);
213 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(
bool vectorized)
const {
214 return TensorOpCost(
sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize);
217 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
218 EvaluatorPointerType data()
const {
return m_buffer; }
222 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void bind(cl::sycl::handler &cgh)
const {
228 TensorEvaluator<ArgType, Device> m_impl;
230 const Device EIGEN_DEVICE_REF m_device;
231 EvaluatorPointerType m_buffer;