// This file is part of Eigen, a lightweight C++ template library // for linear algebra. // // Copyright (C) 2014 Benoit Steiner // // This Source Code Form is subject to the terms of the Mozilla // Public License v. 2.0. If a copy of the MPL was not distributed // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. #ifndef EIGEN_CXX11_TENSOR_TENSOR_FUNCTORS_H #define EIGEN_CXX11_TENSOR_TENSOR_FUNCTORS_H namespace Eigen { namespace internal { /** \internal * \brief Template functor to compute the modulo between an array and a scalar. */ template struct scalar_mod_op { EIGEN_DEVICE_FUNC scalar_mod_op(const Scalar& divisor) : m_divisor(divisor) {} EIGEN_DEVICE_FUNC inline Scalar operator() (const Scalar& a) const { return a % m_divisor; } const Scalar m_divisor; }; template struct functor_traits > { enum { Cost = scalar_div_cost::value, PacketAccess = false }; }; /** \internal * \brief Template functor to compute the modulo between 2 arrays. */ template struct scalar_mod2_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_mod2_op); EIGEN_DEVICE_FUNC inline Scalar operator() (const Scalar& a, const Scalar& b) const { return a % b; } }; template struct functor_traits > { enum { Cost = scalar_div_cost::value, PacketAccess = false }; }; template struct scalar_fmod_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_fmod_op); EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar operator()(const Scalar& a, const Scalar& b) const { return numext::fmod(a, b); } }; template struct functor_traits > { enum { Cost = 13, // Reciprocal throughput of FPREM on Haswell. PacketAccess = false }; }; /** \internal * \brief Template functor to compute the sigmoid of a scalar * \sa class CwiseUnaryOp, ArrayBase::sigmoid() */ template struct scalar_sigmoid_op { EIGEN_EMPTY_STRUCT_CTOR(scalar_sigmoid_op) EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T operator()(const T& x) const { const T one = T(1); return one / (one + numext::exp(-x)); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet packetOp(const Packet& x) const { const Packet one = pset1(T(1)); return pdiv(one, padd(one, pexp(pnegate(x)))); } }; template struct functor_traits > { enum { Cost = NumTraits::AddCost * 2 + NumTraits::MulCost * 6, PacketAccess = packet_traits::HasAdd && packet_traits::HasDiv && packet_traits::HasNegate && packet_traits::HasExp }; }; template struct reducer_traits { enum { Cost = 1, PacketAccess = false }; }; // Standard reduction functors template struct SumReducer { static const bool PacketAccess = packet_traits::HasAdd; static const bool IsStateful = false; EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(const T t, T* accum) const { (*accum) += t; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reducePacket(const Packet& p, Packet* accum) const { (*accum) = padd(*accum, p); } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T initialize() const { internal::scalar_cast_op conv; return conv(0); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet initializePacket() const { return pset1(initialize()); } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalize(const T accum) const { return accum; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet finalizePacket(const Packet& vaccum) const { return vaccum; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalizeBoth(const T saccum, const Packet& vaccum) const { return saccum + predux(vaccum); } }; template struct reducer_traits, Device> { enum { Cost = NumTraits::AddCost, PacketAccess = PacketType::HasAdd }; }; template struct MeanReducer { static const bool PacketAccess = packet_traits::HasAdd && !NumTraits::IsInteger; static const bool IsStateful = true; EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE MeanReducer() : scalarCount_(0), packetCount_(0) { } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(const T t, T* accum) { (*accum) += t; scalarCount_++; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reducePacket(const Packet& p, Packet* accum) { (*accum) = padd(*accum, p); packetCount_++; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T initialize() const { internal::scalar_cast_op conv; return conv(0); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet initializePacket() const { return pset1(initialize()); } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalize(const T accum) const { return accum / scalarCount_; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet finalizePacket(const Packet& vaccum) const { return pdiv(vaccum, pset1(packetCount_)); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalizeBoth(const T saccum, const Packet& vaccum) const { return (saccum + predux(vaccum)) / (scalarCount_ + packetCount_ * unpacket_traits::size); } protected: DenseIndex scalarCount_; DenseIndex packetCount_; }; template struct reducer_traits, Device> { enum { Cost = NumTraits::AddCost, PacketAccess = PacketType::HasAdd }; }; template struct MinMaxBottomValue { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static T bottom_value() { return Eigen::NumTraits::lowest(); } }; template struct MinMaxBottomValue { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static T bottom_value() { return -Eigen::NumTraits::infinity(); } }; template struct MinMaxBottomValue { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static T bottom_value() { return Eigen::NumTraits::highest(); } }; template struct MinMaxBottomValue { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static T bottom_value() { return Eigen::NumTraits::infinity(); } }; template struct MaxReducer { static const bool PacketAccess = packet_traits::HasMax; static const bool IsStateful = false; EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(const T t, T* accum) const { if (t > *accum) { *accum = t; } } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reducePacket(const Packet& p, Packet* accum) const { (*accum) = pmax(*accum, p); } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T initialize() const { return MinMaxBottomValue::IsInteger>::bottom_value(); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet initializePacket() const { return pset1(initialize()); } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalize(const T accum) const { return accum; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet finalizePacket(const Packet& vaccum) const { return vaccum; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalizeBoth(const T saccum, const Packet& vaccum) const { return numext::maxi(saccum, predux_max(vaccum)); } }; template struct reducer_traits, Device> { enum { Cost = NumTraits::AddCost, PacketAccess = PacketType::HasMax }; }; template struct MinReducer { static const bool PacketAccess = packet_traits::HasMin; static const bool IsStateful = false; EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(const T t, T* accum) const { if (t < *accum) { *accum = t; } } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reducePacket(const Packet& p, Packet* accum) const { (*accum) = pmin(*accum, p); } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T initialize() const { return MinMaxBottomValue::IsInteger>::bottom_value(); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet initializePacket() const { return pset1(initialize()); } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalize(const T accum) const { return accum; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet finalizePacket(const Packet& vaccum) const { return vaccum; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalizeBoth(const T saccum, const Packet& vaccum) const { return numext::mini(saccum, predux_min(vaccum)); } }; template struct reducer_traits, Device> { enum { Cost = NumTraits::AddCost, PacketAccess = PacketType::HasMin }; }; template struct ProdReducer { static const bool PacketAccess = packet_traits::HasMul; static const bool IsStateful = false; EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(const T t, T* accum) const { (*accum) *= t; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reducePacket(const Packet& p, Packet* accum) const { (*accum) = pmul(*accum, p); } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T initialize() const { internal::scalar_cast_op conv; return conv(1); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet initializePacket() const { return pset1(initialize()); } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalize(const T accum) const { return accum; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet finalizePacket(const Packet& vaccum) const { return vaccum; } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalizeBoth(const T saccum, const Packet& vaccum) const { return saccum * predux_mul(vaccum); } }; template struct reducer_traits, Device> { enum { Cost = NumTraits::MulCost, PacketAccess = PacketType::HasMul }; }; struct AndReducer { static const bool PacketAccess = false; static const bool IsStateful = false; EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(bool t, bool* accum) const { *accum = *accum && t; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool initialize() const { return true; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool finalize(bool accum) const { return accum; } }; template struct reducer_traits { enum { Cost = 1, PacketAccess = false }; }; struct OrReducer { static const bool PacketAccess = false; static const bool IsStateful = false; EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(bool t, bool* accum) const { *accum = *accum || t; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool initialize() const { return false; } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool finalize(bool accum) const { return accum; } }; template struct reducer_traits { enum { Cost = 1, PacketAccess = false }; }; // Argmin/Argmax reducers template struct ArgMaxTupleReducer { static const bool PacketAccess = false; static const bool IsStateful = false; EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(const T t, T* accum) const { if (t.second > accum->second) { *accum = t; } } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T initialize() const { return T(0, NumTraits::lowest()); } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalize(const T& accum) const { return accum; } }; template struct reducer_traits, Device> { enum { Cost = NumTraits::AddCost, PacketAccess = false }; }; template struct ArgMinTupleReducer { static const bool PacketAccess = false; static const bool IsStateful = false; EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(const T& t, T* accum) const { if (t.second < accum->second) { *accum = t; } } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T initialize() const { return T(0, NumTraits::highest()); } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T finalize(const T& accum) const { return accum; } }; template struct reducer_traits, Device> { enum { Cost = NumTraits::AddCost, PacketAccess = false }; }; // Random number generation namespace { #ifdef __CUDA_ARCH__ __device__ int get_random_seed() { return clock(); } #else static inline int get_random_seed() { #ifdef _WIN32 SYSTEMTIME st; GetSystemTime(&st); return st.wSecond + 1000 * st.wMilliseconds; #elif defined __APPLE__ return static_cast(mach_absolute_time()); #else timespec ts; clock_gettime(CLOCK_REALTIME, &ts); return static_cast(ts.tv_nsec); #endif } #endif } #if !defined (EIGEN_USE_GPU) || !defined(__CUDACC__) || !defined(__CUDA_ARCH__) // We're not compiling a cuda kernel template class UniformRandomGenerator { public: static const bool PacketAccess = true; UniformRandomGenerator(bool deterministic = true) : m_deterministic(deterministic) { if (!deterministic) { srand(get_random_seed()); } } UniformRandomGenerator(const UniformRandomGenerator& other) { m_deterministic = other.m_deterministic; } T operator()() const { return random(); } template PacketType packetOp() const { const int packetSize = internal::unpacket_traits::size; EIGEN_ALIGN_MAX T values[packetSize]; for (int i = 0; i < packetSize; ++i) { values[i] = random(); } return internal::pload(values); } private: bool m_deterministic; }; #if __cplusplus > 199711 || EIGEN_COMP_MSVC >= 1900 template <> class UniformRandomGenerator { public: static const bool PacketAccess = true; UniformRandomGenerator(bool deterministic = true) : m_deterministic(deterministic), m_generator(new std::mt19937()) { if (!deterministic) { m_generator->seed(get_random_seed()); } } UniformRandomGenerator(const UniformRandomGenerator& other) { m_generator = new std::mt19937(); m_generator->seed(other() * UINT_MAX); m_deterministic = other.m_deterministic; } ~UniformRandomGenerator() { delete m_generator; } float operator()() const { return m_distribution(*m_generator); } template PacketType packetOp() const { const int packetSize = internal::unpacket_traits::size; EIGEN_ALIGN_MAX float values[packetSize]; for (int k = 0; k < packetSize; ++k) { values[k] = this->operator()(); } return internal::pload(values); } private: UniformRandomGenerator& operator = (const UniformRandomGenerator&); // Make sure m_deterministic comes first to match the layout of the cpu // version of the code. bool m_deterministic; std::mt19937* m_generator; mutable std::uniform_real_distribution m_distribution; }; template <> class UniformRandomGenerator { public: static const bool PacketAccess = true; UniformRandomGenerator(bool deterministic = true) : m_deterministic(deterministic), m_generator(new std::mt19937()) { if (!deterministic) { m_generator->seed(get_random_seed()); } } UniformRandomGenerator(const UniformRandomGenerator& other) { m_generator = new std::mt19937(); m_generator->seed(other() * UINT_MAX); m_deterministic = other.m_deterministic; } ~UniformRandomGenerator() { delete m_generator; } double operator()() const { return m_distribution(*m_generator); } template PacketType packetOp() const { const int packetSize = internal::unpacket_traits::size; EIGEN_ALIGN_MAX double values[packetSize]; for (int k = 0; k < packetSize; ++k) { values[k] = this->operator()(); } return internal::pload(values); } private: UniformRandomGenerator& operator = (const UniformRandomGenerator&); // Make sure m_deterministic comes first to match the layout of the cpu // version of the code. bool m_deterministic; std::mt19937* m_generator; mutable std::uniform_real_distribution m_distribution; }; #endif #else // We're compiling a cuda kernel template class UniformRandomGenerator; template <> class UniformRandomGenerator { public: static const bool PacketAccess = true; __device__ UniformRandomGenerator(bool deterministic = true) : m_deterministic(deterministic) { const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ UniformRandomGenerator(const UniformRandomGenerator& other) { m_deterministic = other.m_deterministic; const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = m_deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ float operator()() const { return curand_uniform(&m_state); } template __device__ float4 packetOp() const { EIGEN_STATIC_ASSERT((is_same::value), YOU_MADE_A_PROGRAMMING_MISTAKE); return curand_uniform4(&m_state); } private: bool m_deterministic; mutable curandStatePhilox4_32_10_t m_state; }; template <> class UniformRandomGenerator { public: static const bool PacketAccess = true; __device__ UniformRandomGenerator(bool deterministic = true) : m_deterministic(deterministic) { const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ UniformRandomGenerator(const UniformRandomGenerator& other) { m_deterministic = other.m_deterministic; const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = m_deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ double operator()() const { return curand_uniform_double(&m_state); } template __device__ double2 packetOp() const { EIGEN_STATIC_ASSERT((is_same::value), YOU_MADE_A_PROGRAMMING_MISTAKE); return curand_uniform2_double(&m_state); } private: bool m_deterministic; mutable curandStatePhilox4_32_10_t m_state; }; template <> class UniformRandomGenerator > { public: static const bool PacketAccess = false; __device__ UniformRandomGenerator(bool deterministic = true) : m_deterministic(deterministic) { const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ UniformRandomGenerator(const UniformRandomGenerator& other) { m_deterministic = other.m_deterministic; const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = m_deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ std::complex operator()() const { float4 vals = curand_uniform4(&m_state); return std::complex(vals.x, vals.y); } private: bool m_deterministic; mutable curandStatePhilox4_32_10_t m_state; }; template <> class UniformRandomGenerator > { public: static const bool PacketAccess = false; __device__ UniformRandomGenerator(bool deterministic = true) : m_deterministic(deterministic) { const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ UniformRandomGenerator(const UniformRandomGenerator& other) { m_deterministic = other.m_deterministic; const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = m_deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ std::complex operator()() const { double2 vals = curand_uniform2_double(&m_state); return std::complex(vals.x, vals.y); } private: bool m_deterministic; mutable curandStatePhilox4_32_10_t m_state; }; #endif template struct functor_traits > { enum { // Rough estimate. Cost = 100 * NumTraits::MulCost, PacketAccess = UniformRandomGenerator::PacketAccess }; }; #if (!defined (EIGEN_USE_GPU) || !defined(__CUDACC__) || !defined(__CUDA_ARCH__)) && (__cplusplus > 199711 || EIGEN_COMP_MSVC >= 1900) // We're not compiling a cuda kernel template class NormalRandomGenerator { public: static const bool PacketAccess = true; NormalRandomGenerator(bool deterministic = true) : m_deterministic(deterministic), m_distribution(0, 1), m_generator(new std::mt19937()) { if (!deterministic) { m_generator->seed(get_random_seed()); } } NormalRandomGenerator(const NormalRandomGenerator& other) : m_deterministic(other.m_deterministic), m_distribution(other.m_distribution), m_generator(new std::mt19937()) { m_generator->seed(other() * UINT_MAX); } ~NormalRandomGenerator() { delete m_generator; } T operator()() const { return m_distribution(*m_generator); } template PacketType packetOp() const { const int packetSize = internal::unpacket_traits::size; EIGEN_ALIGN_MAX T values[packetSize]; for (int i = 0; i < packetSize; ++i) { values[i] = m_distribution(*m_generator); } return internal::pload(values); } private: // No assignment NormalRandomGenerator& operator = (const NormalRandomGenerator&); bool m_deterministic; mutable std::normal_distribution m_distribution; std::mt19937* m_generator; }; #elif defined (EIGEN_USE_GPU) && defined(__CUDACC__) && defined(__CUDA_ARCH__) // We're compiling a cuda kernel template class NormalRandomGenerator; template <> class NormalRandomGenerator { public: static const bool PacketAccess = true; __device__ NormalRandomGenerator(bool deterministic = true) : m_deterministic(deterministic) { const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ NormalRandomGenerator(const NormalRandomGenerator& other) { m_deterministic = other.m_deterministic; const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = m_deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ float operator()() const { return curand_normal(&m_state); } template __device__ float4 packetOp() const { EIGEN_STATIC_ASSERT((is_same::value), YOU_MADE_A_PROGRAMMING_MISTAKE); return curand_normal4(&m_state); } private: bool m_deterministic; mutable curandStatePhilox4_32_10_t m_state; }; template <> class NormalRandomGenerator { public: static const bool PacketAccess = true; __device__ NormalRandomGenerator(bool deterministic = true) : m_deterministic(deterministic) { const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ NormalRandomGenerator(const NormalRandomGenerator& other) { m_deterministic = other.m_deterministic; const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = m_deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ double operator()() const { return curand_normal_double(&m_state); } template __device__ double2 packetOp() const { EIGEN_STATIC_ASSERT((is_same::value), YOU_MADE_A_PROGRAMMING_MISTAKE); return curand_normal2_double(&m_state); } private: bool m_deterministic; mutable curandStatePhilox4_32_10_t m_state; }; template <> class NormalRandomGenerator > { public: static const bool PacketAccess = false; __device__ NormalRandomGenerator(bool deterministic = true) : m_deterministic(deterministic) { const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ NormalRandomGenerator(const NormalRandomGenerator& other) { m_deterministic = other.m_deterministic; const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = m_deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ std::complex operator()() const { float4 vals = curand_normal4(&m_state); return std::complex(vals.x, vals.y); } private: bool m_deterministic; mutable curandStatePhilox4_32_10_t m_state; }; template <> class NormalRandomGenerator > { public: static const bool PacketAccess = false; __device__ NormalRandomGenerator(bool deterministic = true) : m_deterministic(deterministic) { const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ NormalRandomGenerator(const NormalRandomGenerator& other) { m_deterministic = other.m_deterministic; const int tid = blockIdx.x * blockDim.x + threadIdx.x; const int seed = m_deterministic ? 0 : get_random_seed(); curand_init(seed, tid, 0, &m_state); } __device__ std::complex operator()() const { double2 vals = curand_normal2_double(&m_state); return std::complex(vals.x, vals.y); } private: bool m_deterministic; mutable curandStatePhilox4_32_10_t m_state; }; #else template class NormalRandomGenerator { public: static const bool PacketAccess = false; NormalRandomGenerator(bool deterministic = true) : m_deterministic(deterministic) {} private: bool m_deterministic; }; #endif template struct functor_traits > { enum { // Rough estimate. Cost = 100 * NumTraits::MulCost, PacketAccess = NormalRandomGenerator::PacketAccess }; }; template class GaussianGenerator { public: static const bool PacketAccess = false; EIGEN_DEVICE_FUNC GaussianGenerator(const array& means, const array& std_devs) : m_means(means) { for (size_t i = 0; i < NumDims; ++i) { m_two_sigmas[i] = std_devs[i] * std_devs[i] * 2; } } T operator()(const array& coordinates) const { T tmp = T(0); for (size_t i = 0; i < NumDims; ++i) { T offset = coordinates[i] - m_means[i]; tmp += offset * offset / m_two_sigmas[i]; } return numext::exp(-tmp); } private: array m_means; array m_two_sigmas; }; template struct functor_traits > { enum { Cost = NumDims * (2 * NumTraits::AddCost + NumTraits::MulCost + functor_traits >::Cost) + functor_traits >::Cost, PacketAccess = GaussianGenerator::PacketAccess }; }; } // end namespace internal } // end namespace Eigen #endif // EIGEN_CXX11_TENSOR_TENSOR_FUNCTORS_H