#ifndef EIGEN_CXX11_TENSOR_TENSOR_EVALUATOR_H
#define EIGEN_CXX11_TENSOR_TENSOR_EVALUATOR_H
#include "./InternalHeaderCheck.h"
namespace Eigen {
template <typename Derived, typename Device>
struct TensorEvaluator { … };
namespace internal {
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T loadConstant(const T* address) { … }
#if defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 350
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float loadConstant(const float* address) {
return __ldg(address);
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double loadConstant(const double* address) {
return __ldg(address);
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half loadConstant(const Eigen::half* address) {
return Eigen::half(half_impl::raw_uint16_to_half(__ldg(&address->x)));
}
#endif
}
TensorEvaluator<const Derived, Device>;
TensorEvaluator<const TensorCwiseNullaryOp<NullaryOp, ArgType>, Device>;
TensorEvaluator<const TensorCwiseUnaryOp<UnaryOp, ArgType>, Device>;
TensorEvaluator<const TensorCwiseBinaryOp<BinaryOp, LeftArgType, RightArgType>, Device>;
TensorEvaluator<const TensorCwiseTernaryOp<TernaryOp, Arg1Type, Arg2Type, Arg3Type>, Device>;
TensorEvaluator<const TensorSelectOp<IfArgType, ThenArgType, ElseArgType>, Device>;
}
#if defined(EIGEN_USE_SYCL) && defined(SYCL_COMPILER_IS_DPCPP)
template <typename Derived, typename Device>
struct cl::sycl::is_device_copyable<
Eigen::TensorEvaluator<Derived, Device>,
std::enable_if_t<!std::is_trivially_copyable<Eigen::TensorEvaluator<Derived, Device>>::value>> : std::true_type {};
#endif
#endif