Fixed the tensor shuffling test
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDimensions.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDimensions.h index 11590b4..732c6b3 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDimensions.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDimensions.h
@@ -37,8 +37,7 @@ Index second; }; - -// Boiler plate code +// Boilerplate code namespace internal { template<std::size_t n, typename Dimension> struct dget { @@ -110,6 +109,11 @@ } }; +template <typename std::size_t... Indices> +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::size_t array_prod(const Sizes<Indices...>&) { + return Sizes<Indices...>::total_size; +} + #else template <std::size_t n> @@ -136,9 +140,21 @@ // todo: add assertion } #ifdef EIGEN_HAS_VARIADIC_TEMPLATES + template <typename... DenseIndex> Sizes(DenseIndex... indices) { } explicit Sizes(std::initializer_list<std::size_t> l) { // todo: add assertion } +#else + EIGEN_DEVICE_FUNC explicit Sizes(const DenseIndex i0) { + } + EIGEN_DEVICE_FUNC explicit Sizes(const DenseIndex i0, const DenseIndex i1) { + } + EIGEN_DEVICE_FUNC explicit Sizes(const DenseIndex i0, const DenseIndex i1, const DenseIndex i2) { + } + EIGEN_DEVICE_FUNC explicit Sizes(const DenseIndex i0, const DenseIndex i1, const DenseIndex i2, const DenseIndex i3) { + } + EIGEN_DEVICE_FUNC explicit Sizes(const DenseIndex i0, const DenseIndex i1, const DenseIndex i2, const DenseIndex i3, const DenseIndex i4) { + } #endif template <typename T> Sizes& operator = (const T& other) { @@ -156,9 +172,14 @@ } }; +template <std::size_t V1, std::size_t V2, std::size_t V3, std::size_t V4, std::size_t V5> +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::size_t array_prod(const Sizes<V1, V2, V3, V4, V5>&) { + return Sizes<V1, V2, V3, V4, V5>::total_size; +}; + #endif -// Boiler plate +// Boilerplate namespace internal { template<typename Index, std::size_t NumIndices, std::size_t n, bool RowMajor> struct tensor_index_linearization_helper @@ -243,6 +264,112 @@ }; + + +// Boilerplate +namespace internal { +template<typename Index, std::size_t NumIndices, std::size_t n, bool RowMajor> +struct tensor_vsize_index_linearization_helper +{ + static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE + Index run(array<Index, NumIndices> const& indices, std::vector<DenseIndex> const& dimensions) + { + return array_get<RowMajor ? n : (NumIndices - n - 1)>(indices) + + array_get<RowMajor ? n : (NumIndices - n - 1)>(dimensions) * + tensor_vsize_index_linearization_helper<Index, NumIndices, n - 1, RowMajor>::run(indices, dimensions); + } +}; + +template<typename Index, std::size_t NumIndices, bool RowMajor> +struct tensor_vsize_index_linearization_helper<Index, NumIndices, 0, RowMajor> +{ + static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE + Index run(array<Index, NumIndices> const& indices, std::vector<DenseIndex> const&) + { + return array_get<RowMajor ? 0 : NumIndices - 1>(indices); + } +}; +} // end namespace internal + +template <typename DenseIndex> +struct VSizes : std::vector<DenseIndex> { + typedef std::vector<DenseIndex> Base; + + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE size_t TotalSize() const { + return internal::array_prod(*static_cast<const Base*>(this)); + } + + EIGEN_DEVICE_FUNC VSizes() { } + EIGEN_DEVICE_FUNC explicit VSizes(const std::vector<DenseIndex>& a) : Base(a) { } + + template <std::size_t NumDims> + EIGEN_DEVICE_FUNC explicit VSizes(const array<DenseIndex, NumDims>& a) { + this->resize(NumDims); + for (int i = 0; i < NumDims; ++i) { + (*this)[i] = a[i]; + } + } + + EIGEN_DEVICE_FUNC explicit VSizes(const DenseIndex i0) { + this->resize(1); + (*this)[0] = i0; + } + EIGEN_DEVICE_FUNC explicit VSizes(const DenseIndex i0, const DenseIndex i1) { + this->resize(2); + (*this)[0] = i0; + (*this)[1] = i1; + } + EIGEN_DEVICE_FUNC explicit VSizes(const DenseIndex i0, const DenseIndex i1, const DenseIndex i2) { + this->resize(3); + (*this)[0] = i0; + (*this)[1] = i1; + (*this)[2] = i2; + } + EIGEN_DEVICE_FUNC explicit VSizes(const DenseIndex i0, const DenseIndex i1, const DenseIndex i2, const DenseIndex i3) { + this->resize(4); + (*this)[0] = i0; + (*this)[1] = i1; + (*this)[2] = i2; + (*this)[3] = i3; + } + EIGEN_DEVICE_FUNC explicit VSizes(const DenseIndex i0, const DenseIndex i1, const DenseIndex i2, const DenseIndex i3, const DenseIndex i4) { + this->resize(5); + (*this)[0] = i0; + (*this)[1] = i1; + (*this)[2] = i2; + (*this)[3] = i3; + (*this)[4] = i4; + } + + VSizes& operator = (const std::vector<DenseIndex>& other) { + *static_cast<Base*>(this) = other; + return *this; + } + + // A constexpr would be so much better here + template <std::size_t NumDims> + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE size_t IndexOfColMajor(const array<DenseIndex, NumDims>& indices) const { + return internal::tensor_vsize_index_linearization_helper<DenseIndex, NumDims, NumDims - 1, false>::run(indices, *static_cast<const Base*>(this)); + } + template <std::size_t NumDims> + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE size_t IndexOfRowMajor(const array<DenseIndex, NumDims>& indices) const { + return internal::tensor_vsize_index_linearization_helper<DenseIndex, NumDims, NumDims - 1, true>::run(indices, *static_cast<const Base*>(this)); + } +}; + + +// Boilerplate +namespace internal { +template <typename DenseIndex> +EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE DenseIndex array_prod(const VSizes<DenseIndex>& sizes) { + DenseIndex total_size = 1; + for (int i = 0; i < sizes.size(); ++i) { + total_size *= sizes[i]; + } + return total_size; +} +} + namespace internal { template <typename DenseIndex, std::size_t NumDims> struct array_size<const DSizes<DenseIndex, NumDims> > {
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h index 84768ca..10f5a5e 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h
@@ -39,7 +39,7 @@ const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL); if (needs_assign) { - const Index size = evaluator.dimensions().TotalSize(); + const Index size = array_prod(evaluator.dimensions()); for (Index i = 0; i < size; ++i) { evaluator.evalScalar(i); } @@ -60,7 +60,7 @@ const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL); if (needs_assign) { - const Index size = evaluator.dimensions().TotalSize(); + const Index size = array_prod(evaluator.dimensions()); static const int PacketSize = unpacket_traits<typename TensorEvaluator<Expression, DefaultDevice>::PacketReturnType>::size; const int VectorizedSize = (size / PacketSize) * PacketSize; @@ -122,7 +122,7 @@ const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL); if (needs_assign) { - const Index size = evaluator.dimensions().TotalSize(); + const Index size = array_prod(evaluator.dimensions()); static const int PacketSize = Vectorizable ? unpacket_traits<typename Evaluator::PacketReturnType>::size : 1; @@ -176,7 +176,7 @@ const int num_blocks = getNumCudaMultiProcessors() * maxCudaThreadsPerMultiProcessor() / maxCudaThreadsPerBlock(); const int block_size = maxCudaThreadsPerBlock(); - const Index size = evaluator.dimensions().TotalSize(); + const Index size = array_prod(evaluator.dimensions()); EigenMetaKernel<TensorEvaluator<Expression, GpuDevice> > <<<num_blocks, block_size, 0, device.stream()>>>(evaluator, size); assert(cudaGetLastError() == cudaSuccess); }
diff --git a/unsupported/test/CMakeLists.txt b/unsupported/test/CMakeLists.txt index 1c4d083..ac2ccaf 100644 --- a/unsupported/test/CMakeLists.txt +++ b/unsupported/test/CMakeLists.txt
@@ -119,7 +119,7 @@ ei_add_test(cxx11_tensor_morphing "-std=c++0x") ei_add_test(cxx11_tensor_padding "-std=c++0x") ei_add_test(cxx11_tensor_reduction "-std=c++0x") -# ei_add_test(cxx11_tensor_shuffling "-std=c++0x") + ei_add_test(cxx11_tensor_shuffling "-std=c++0x") ei_add_test(cxx11_tensor_striding "-std=c++0x") # ei_add_test(cxx11_tensor_device "-std=c++0x") ei_add_test(cxx11_tensor_thread_pool "-std=c++0x")
diff --git a/unsupported/test/cxx11_tensor_fixed_size.cpp b/unsupported/test/cxx11_tensor_fixed_size.cpp index d270486..b0501aa 100644 --- a/unsupported/test/cxx11_tensor_fixed_size.cpp +++ b/unsupported/test/cxx11_tensor_fixed_size.cpp
@@ -179,7 +179,7 @@ for (int i = 0; i < 2; ++i) { for (int j = 0; j < 3; ++j) { for (int k = 0; k < 7; ++k) { - VERIFY_IS_APPROX(mat3(array<ptrdiff_t, 3>(i,j,k)), powf(val, 3.5f)); + VERIFY_IS_APPROX(mat3(i,j,k), powf(val, 3.5f)); val += 1.0; } }
diff --git a/unsupported/test/cxx11_tensor_shuffling.cpp b/unsupported/test/cxx11_tensor_shuffling.cpp index 5ab8b68..39c6234 100644 --- a/unsupported/test/cxx11_tensor_shuffling.cpp +++ b/unsupported/test/cxx11_tensor_shuffling.cpp
@@ -12,6 +12,7 @@ #include <Eigen/CXX11/Tensor> using Eigen::Tensor; +using Eigen::array; static void test_simple_shuffling() { @@ -80,10 +81,10 @@ Tensor<float, 4> result(5,7,3,2); - array<int, 4> src_slice_dim(Eigen::array<int, 4>(2,3,1,7)); - array<int, 4> src_slice_start(Eigen::array<int, 4>(0,0,0,0)); - array<int, 4> dst_slice_dim(Eigen::array<int, 4>(1,7,3,2)); - array<int, 4> dst_slice_start(Eigen::array<int, 4>(0,0,0,0)); + array<int, 4> src_slice_dim{{2,3,1,7}}; + array<int, 4> src_slice_start{{0,0,0,0}}; + array<int, 4> dst_slice_dim{{1,7,3,2}}; + array<int, 4> dst_slice_start{{0,0,0,0}}; for (int i = 0; i < 5; ++i) { result.slice(dst_slice_start, dst_slice_dim) =