10#ifndef EIGEN_CXX11_TENSOR_TENSOR_CUSTOM_OP_H
11#define EIGEN_CXX11_TENSOR_TENSOR_CUSTOM_OP_H
23template<
typename CustomUnaryFunc,
typename XprType>
24struct traits<TensorCustomUnaryOp<CustomUnaryFunc, XprType> >
26 typedef typename XprType::Scalar Scalar;
27 typedef typename XprType::StorageKind StorageKind;
28 typedef typename XprType::Index
Index;
29 typedef typename XprType::Nested Nested;
30 typedef typename remove_reference<Nested>::type _Nested;
31 static const int NumDimensions = traits<XprType>::NumDimensions;
32 static const int Layout = traits<XprType>::Layout;
33 typedef typename traits<XprType>::PointerType PointerType;
36template<
typename CustomUnaryFunc,
typename XprType>
37struct eval<TensorCustomUnaryOp<CustomUnaryFunc, XprType>,
Eigen::Dense>
39 typedef const TensorCustomUnaryOp<CustomUnaryFunc, XprType>EIGEN_DEVICE_REF type;
42template<
typename CustomUnaryFunc,
typename XprType>
43struct nested<TensorCustomUnaryOp<CustomUnaryFunc, XprType> >
45 typedef TensorCustomUnaryOp<CustomUnaryFunc, XprType> type;
52template<
typename CustomUnaryFunc,
typename XprType>
56 typedef typename internal::traits<TensorCustomUnaryOp>::Scalar Scalar;
58 typedef typename XprType::CoeffReturnType CoeffReturnType;
59 typedef typename internal::nested<TensorCustomUnaryOp>::type Nested;
60 typedef typename internal::traits<TensorCustomUnaryOp>::StorageKind StorageKind;
61 typedef typename internal::traits<TensorCustomUnaryOp>::Index Index;
63 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
TensorCustomUnaryOp(
const XprType& expr,
const CustomUnaryFunc& func)
64 : m_expr(expr), m_func(func) {}
67 const CustomUnaryFunc& func()
const {
return m_func; }
70 const typename internal::remove_all<typename XprType::Nested>::type&
71 expression()
const {
return m_expr; }
74 typename XprType::Nested m_expr;
75 const CustomUnaryFunc m_func;
80template<
typename CustomUnaryFunc,
typename XprType,
typename Device>
84 typedef typename internal::traits<ArgType>::Index
Index;
85 static const int NumDims = internal::traits<ArgType>::NumDimensions;
86 typedef DSizes<Index, NumDims> Dimensions;
87 typedef typename internal::remove_const<typename ArgType::Scalar>::type Scalar;
88 typedef typename internal::remove_const<typename XprType::CoeffReturnType>::type CoeffReturnType;
89 typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
90 static const int PacketSize = PacketType<CoeffReturnType, Device>::size;
91 typedef typename Eigen::internal::traits<XprType>::PointerType TensorPointerType;
92 typedef StorageMemory<CoeffReturnType, Device> Storage;
93 typedef typename Storage::Type EvaluatorPointerType;
97 PacketAccess = (PacketType<CoeffReturnType, Device>::size > 1),
99 PreferBlockAccess =
false,
106 typedef internal::TensorBlockNotImplemented TensorBlock;
109 EIGEN_STRONG_INLINE TensorEvaluator(
const ArgType& op,
const Device& device)
110 : m_op(op), m_device(device), m_result(NULL)
112 m_dimensions = op.func().dimensions(op.expression());
115 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
const Dimensions& dimensions()
const {
return m_dimensions; }
117 EIGEN_STRONG_INLINE
bool evalSubExprsIfNeeded(EvaluatorPointerType data) {
122 m_result =
static_cast<EvaluatorPointerType
>(m_device.get( (CoeffReturnType*)
123 m_device.allocate_temp(dimensions().TotalSize() *
sizeof(Scalar))));
129 EIGEN_STRONG_INLINE
void cleanup() {
131 m_device.deallocate_temp(m_result);
136 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index)
const {
137 return m_result[index];
140 template<
int LoadMode>
141 EIGEN_DEVICE_FUNC PacketReturnType packet(Index index)
const {
142 return internal::ploadt<PacketReturnType, LoadMode>(m_result + index);
145 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(
bool vectorized)
const {
147 return TensorOpCost(
sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize);
150 EIGEN_DEVICE_FUNC EvaluatorPointerType data()
const {
return m_result; }
154 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void bind(cl::sycl::handler &cgh)
const {
160 void evalTo(EvaluatorPointerType data) {
161 TensorMap<Tensor<CoeffReturnType, NumDims, Layout, Index> > result(m_device.get(data), m_dimensions);
162 m_op.func().eval(m_op.expression(), result, m_device);
165 Dimensions m_dimensions;
167 const Device EIGEN_DEVICE_REF m_device;
168 EvaluatorPointerType m_result;
181template<
typename CustomBinaryFunc,
typename LhsXprType,
typename RhsXprType>
182struct traits<TensorCustomBinaryOp<CustomBinaryFunc, LhsXprType, RhsXprType> >
184 typedef typename internal::promote_storage_type<
typename LhsXprType::Scalar,
185 typename RhsXprType::Scalar>::ret Scalar;
186 typedef typename internal::promote_storage_type<
typename LhsXprType::CoeffReturnType,
187 typename RhsXprType::CoeffReturnType>::ret CoeffReturnType;
188 typedef typename promote_storage_type<typename traits<LhsXprType>::StorageKind,
189 typename traits<RhsXprType>::StorageKind>::ret StorageKind;
190 typedef typename promote_index_type<typename traits<LhsXprType>::Index,
191 typename traits<RhsXprType>::Index>::type
Index;
192 typedef typename LhsXprType::Nested LhsNested;
193 typedef typename RhsXprType::Nested RhsNested;
194 typedef typename remove_reference<LhsNested>::type _LhsNested;
195 typedef typename remove_reference<RhsNested>::type _RhsNested;
196 static const int NumDimensions = traits<LhsXprType>::NumDimensions;
197 static const int Layout = traits<LhsXprType>::Layout;
198 typedef typename conditional<Pointer_type_promotion<typename LhsXprType::Scalar, Scalar>::val,
199 typename traits<LhsXprType>::PointerType,
typename traits<RhsXprType>::PointerType>::type PointerType;
202template<
typename CustomBinaryFunc,
typename LhsXprType,
typename RhsXprType>
203struct eval<TensorCustomBinaryOp<CustomBinaryFunc, LhsXprType, RhsXprType>,
Eigen::Dense>
205 typedef const TensorCustomBinaryOp<CustomBinaryFunc, LhsXprType, RhsXprType>& type;
208template<
typename CustomBinaryFunc,
typename LhsXprType,
typename RhsXprType>
209struct nested<TensorCustomBinaryOp<CustomBinaryFunc, LhsXprType, RhsXprType> >
211 typedef TensorCustomBinaryOp<CustomBinaryFunc, LhsXprType, RhsXprType> type;
218template<
typename CustomBinaryFunc,
typename LhsXprType,
typename RhsXprType>
222 typedef typename internal::traits<TensorCustomBinaryOp>::Scalar Scalar;
224 typedef typename internal::traits<TensorCustomBinaryOp>::CoeffReturnType CoeffReturnType;
225 typedef typename internal::nested<TensorCustomBinaryOp>::type Nested;
226 typedef typename internal::traits<TensorCustomBinaryOp>::StorageKind StorageKind;
227 typedef typename internal::traits<TensorCustomBinaryOp>::Index Index;
229 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
TensorCustomBinaryOp(
const LhsXprType& lhs,
const RhsXprType& rhs,
const CustomBinaryFunc& func)
231 : m_lhs_xpr(lhs), m_rhs_xpr(rhs), m_func(func) {}
234 const CustomBinaryFunc& func()
const {
return m_func; }
237 const typename internal::remove_all<typename LhsXprType::Nested>::type&
238 lhsExpression()
const {
return m_lhs_xpr; }
241 const typename internal::remove_all<typename RhsXprType::Nested>::type&
242 rhsExpression()
const {
return m_rhs_xpr; }
245 typename LhsXprType::Nested m_lhs_xpr;
246 typename RhsXprType::Nested m_rhs_xpr;
247 const CustomBinaryFunc m_func;
252template<
typename CustomBinaryFunc,
typename LhsXprType,
typename RhsXprType,
typename Device>
256 typedef typename internal::traits<XprType>::Index
Index;
257 static const int NumDims = internal::traits<XprType>::NumDimensions;
258 typedef DSizes<Index, NumDims> Dimensions;
259 typedef typename XprType::Scalar Scalar;
260 typedef typename internal::remove_const<typename XprType::CoeffReturnType>::type CoeffReturnType;
261 typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
262 static const int PacketSize = PacketType<CoeffReturnType, Device>::size;
264 typedef typename Eigen::internal::traits<XprType>::PointerType TensorPointerType;
265 typedef StorageMemory<CoeffReturnType, Device> Storage;
266 typedef typename Storage::Type EvaluatorPointerType;
270 PacketAccess = (PacketType<CoeffReturnType, Device>::size > 1),
272 PreferBlockAccess =
false,
279 typedef internal::TensorBlockNotImplemented TensorBlock;
282 EIGEN_STRONG_INLINE TensorEvaluator(
const XprType& op,
const Device& device)
283 : m_op(op), m_device(device), m_result(NULL)
285 m_dimensions = op.func().dimensions(op.lhsExpression(), op.rhsExpression());
288 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
const Dimensions& dimensions()
const {
return m_dimensions; }
290 EIGEN_STRONG_INLINE
bool evalSubExprsIfNeeded(EvaluatorPointerType data) {
295 m_result =
static_cast<EvaluatorPointerType
>(m_device.get( (CoeffReturnType*)
296 m_device.allocate_temp(dimensions().TotalSize() *
sizeof(CoeffReturnType))));
302 EIGEN_STRONG_INLINE
void cleanup() {
303 if (m_result != NULL) {
304 m_device.deallocate_temp(m_result);
309 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE CoeffReturnType coeff(Index index)
const {
310 return m_result[index];
313 template<
int LoadMode>
314 EIGEN_DEVICE_FUNC PacketReturnType packet(Index index)
const {
315 return internal::ploadt<PacketReturnType, LoadMode>(m_result + index);
318 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorOpCost costPerCoeff(
bool vectorized)
const {
320 return TensorOpCost(
sizeof(CoeffReturnType), 0, 0, vectorized, PacketSize);
323 EIGEN_DEVICE_FUNC EvaluatorPointerType data()
const {
return m_result; }
327 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void bind(cl::sycl::handler &cgh)
const {
333 void evalTo(EvaluatorPointerType data) {
334 TensorMap<Tensor<CoeffReturnType, NumDims, Layout> > result(m_device.get(data), m_dimensions);
335 m_op.func().eval(m_op.lhsExpression(), m_op.rhsExpression(), result, m_device);
338 Dimensions m_dimensions;
340 const Device EIGEN_DEVICE_REF m_device;
341 EvaluatorPointerType m_result;
The tensor base class.
Definition: TensorForwardDeclarations.h:56
Tensor custom class.
Definition: TensorCustomOp.h:220
Tensor custom class.
Definition: TensorCustomOp.h:54
Namespace containing all symbols from the Eigen library.
EIGEN_DEFAULT_DENSE_INDEX_TYPE Index
A cost model used to limit the number of threads used for evaluating tensor expression.
Definition: TensorEvaluator.h:29