diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorFunctors.h b/unsupported/Eigen/CXX11/src/Tensor/TensorFunctors.h index 92984336c..e9aa22183 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorFunctors.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorFunctors.h @@ -25,12 +25,12 @@ template struct SumReducer } private: - T m_sum; + typename internal::remove_all::type m_sum; }; template struct MaxReducer { - EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE MaxReducer() : m_max((std::numeric_limits::min)()) { } + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE MaxReducer() : m_max(-(std::numeric_limits::max)()) { } EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void reduce(const T t) { if (t > m_max) { m_max = t; } } @@ -39,7 +39,7 @@ template struct MaxReducer } private: - T m_max; + typename internal::remove_all::type m_max; }; template struct MinReducer @@ -53,9 +53,73 @@ template struct MinReducer } private: - T m_min; + typename internal::remove_all::type m_min; }; + +#if !defined (EIGEN_USE_GPU) || !defined(__CUDACC__) || !defined(__CUDA_ARCH__) +// We're not compiling a cuda kernel +template struct UniformRandomGenerator { + template + T operator()(Index, Index = 0) const { + return random(); + } + template + typename internal::packet_traits::type packetOp(Index, Index = 0) const { + const int packetSize = internal::packet_traits::size; + EIGEN_ALIGN_DEFAULT T values[packetSize]; + for (int i = 0; i < packetSize; ++i) { + values[i] = random(); + } + return internal::pload::type>(values); + } +}; + +#else + +// We're compiling a cuda kernel +template struct UniformRandomGenerator; + +template <> struct UniformRandomGenerator { + UniformRandomGenerator() { + const int tid = blockIdx.x * blockDim.x + threadIdx.x; + curand_init(0, tid, 0, &m_state); + } + + template + float operator()(Index, Index = 0) const { + return curand_uniform(&m_state); + } + template + float4 packetOp(Index, Index = 0) const { + return curand_uniform4(&m_state); + } + + private: + mutable curandStatePhilox4_32_10_t m_state; +}; + +template <> struct UniformRandomGenerator { + UniformRandomGenerator() { + const int tid = blockIdx.x * blockDim.x + threadIdx.x; + curand_init(0, tid, 0, &m_state); + } + template + double operator()(Index, Index = 0) const { + return curand_uniform_double(&m_state); + } + template + double2 packetOp(Index, Index = 0) const { + return curand_uniform2_double(&m_state); + } + + private: + mutable curandStatePhilox4_32_10_t m_state; +}; + +#endif + + } // end namespace internal } // end namespace Eigen diff --git a/unsupported/test/cxx11_tensor_reduction.cpp b/unsupported/test/cxx11_tensor_reduction.cpp index 27135b982..da9885166 100644 --- a/unsupported/test/cxx11_tensor_reduction.cpp +++ b/unsupported/test/cxx11_tensor_reduction.cpp @@ -139,9 +139,42 @@ static void test_user_defined_reductions() } +static void test_tensor_maps() +{ + int inputs[2*3*5*7]; + TensorMap > tensor_map(inputs, 2,3,5,7); + TensorMap > tensor_map_const(inputs, 2,3,5,7); + const TensorMap > tensor_map_const_const(inputs, 2,3,5,7); + + tensor_map.setRandom(); + array reduction_axis; + reduction_axis[0] = 1; + reduction_axis[1] = 3; + + Tensor result = tensor_map.sum(reduction_axis); + Tensor result2 = tensor_map_const.sum(reduction_axis); + Tensor result3 = tensor_map_const_const.sum(reduction_axis); + + for (int i = 0; i < 2; ++i) { + for (int j = 0; j < 5; ++j) { + int sum = 0; + for (int k = 0; k < 3; ++k) { + for (int l = 0; l < 7; ++l) { + sum += tensor_map(i, k, j, l); + } + } + VERIFY_IS_EQUAL(result(i, j), sum); + VERIFY_IS_EQUAL(result2(i, j), sum); + VERIFY_IS_EQUAL(result3(i, j), sum); + } + } +} + + void test_cxx11_tensor_reduction() { CALL_SUBTEST(test_simple_reductions()); CALL_SUBTEST(test_full_reductions()); CALL_SUBTEST(test_user_defined_reductions()); + CALL_SUBTEST(test_tensor_maps()); }