diff options
| author |  Benoit Steiner <benoit.steiner.goog@gmail.com> | 2015-01-14 15:46:04 -0800 |
|---|---|---|
| committer | Benoit Steiner <benoit.steiner.goog@gmail.com> | 2015-01-14 15:46:04 -0800 |
| commit | b5124e7cfda27ed99dcfcec8cb1b674efa1ef4a3 (patch) | |
| tree | 7f8378843a756af14785e563689b4765e062a953 | |
| parent | 54e3633b437e44ed4d370c9f8868535192308ca3 (diff) | |
Created many additional tests
23 files changed, 1908 insertions, 367 deletions
diff --git a/unsupported/test/CMakeLists.txt b/unsupported/test/CMakeLists.txt index 89c651804..9f44e47f9 100644 --- a/unsupported/test/CMakeLists.txt +++ b/unsupported/test/CMakeLists.txt @@ -99,7 +99,7 @@ if(EIGEN_TEST_CXX11) # older compiler that don't support cxx11. ei_add_test(cxx11_meta "-std=c++0x") ei_add_test(cxx11_tensor_simple "-std=c++0x") - ei_add_test(cxx11_tensor_symmetry "-std=c++0x") +# ei_add_test(cxx11_tensor_symmetry "-std=c++0x") ei_add_test(cxx11_tensor_assign "-std=c++0x") ei_add_test(cxx11_tensor_dimension "-std=c++0x") ei_add_test(cxx11_tensor_index_list "-std=c++0x") @@ -126,8 +126,17 @@ if(EIGEN_TEST_CXX11) ei_add_test(cxx11_tensor_reduction "-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") ei_add_test(cxx11_tensor_ref "-std=c++0x") + ei_add_test(cxx11_tensor_random "-std=c++0x") + ei_add_test(cxx11_tensor_casts "-std=c++0x") + ei_add_test(cxx11_tensor_reverse "-std=c++0x") + ei_add_test(cxx11_tensor_layout_swap "-std=c++0x") ei_add_test(cxx11_tensor_io "-std=c++0x") + + # These tests needs nvcc +# ei_add_test(cxx11_tensor_device "-std=c++0x") +# ei_add_test(cxx11_tensor_cuda "-std=c++0x") +# ei_add_test(cxx11_tensor_contract_cuda "-std=c++0x") + endif() diff --git a/unsupported/test/cxx11_tensor_assign.cpp b/unsupported/test/cxx11_tensor_assign.cpp index 0ac3f9bf9..d16aaf847 100644 --- a/unsupported/test/cxx11_tensor_assign.cpp +++ b/unsupported/test/cxx11_tensor_assign.cpp @@ -285,6 +285,78 @@ static void test_compound_assign() } } +static void test_std_initializers_tensor() { +#ifdef EIGEN_HAS_VARIADIC_TEMPLATES + Tensor<int, 1> a(3); + a.setValues({0, 1, 2}); + VERIFY_IS_EQUAL(a(0), 0); + VERIFY_IS_EQUAL(a(1), 1); + VERIFY_IS_EQUAL(a(2), 2); + + // It fills the top-left slice. + a.setValues({10, 20}); + VERIFY_IS_EQUAL(a(0), 10); + VERIFY_IS_EQUAL(a(1), 20); + VERIFY_IS_EQUAL(a(2), 2); + + // Chaining. + Tensor<int, 1> a2(3); + a2 = a.setValues({100, 200, 300}); + VERIFY_IS_EQUAL(a(0), 100); + VERIFY_IS_EQUAL(a(1), 200); + VERIFY_IS_EQUAL(a(2), 300); + VERIFY_IS_EQUAL(a2(0), 100); + VERIFY_IS_EQUAL(a2(1), 200); + VERIFY_IS_EQUAL(a2(2), 300); + + Tensor<int, 2> b(2, 3); + b.setValues({{0, 1, 2}, {3, 4, 5}}); + VERIFY_IS_EQUAL(b(0, 0), 0); + VERIFY_IS_EQUAL(b(0, 1), 1); + VERIFY_IS_EQUAL(b(0, 2), 2); + VERIFY_IS_EQUAL(b(1, 0), 3); + VERIFY_IS_EQUAL(b(1, 1), 4); + VERIFY_IS_EQUAL(b(1, 2), 5); + + // It fills the top-left slice. + b.setValues({{10, 20}, {30}}); + VERIFY_IS_EQUAL(b(0, 0), 10); + VERIFY_IS_EQUAL(b(0, 1), 20); + VERIFY_IS_EQUAL(b(0, 2), 2); + VERIFY_IS_EQUAL(b(1, 0), 30); + VERIFY_IS_EQUAL(b(1, 1), 4); + VERIFY_IS_EQUAL(b(1, 2), 5); + + Eigen::Tensor<int, 3> c(3, 2, 4); + c.setValues({{{0, 1, 2, 3}, {4, 5, 6, 7}}, + {{10, 11, 12, 13}, {14, 15, 16, 17}}, + {{20, 21, 22, 23}, {24, 25, 26, 27}}}); + VERIFY_IS_EQUAL(c(0, 0, 0), 0); + VERIFY_IS_EQUAL(c(0, 0, 1), 1); + VERIFY_IS_EQUAL(c(0, 0, 2), 2); + VERIFY_IS_EQUAL(c(0, 0, 3), 3); + VERIFY_IS_EQUAL(c(0, 1, 0), 4); + VERIFY_IS_EQUAL(c(0, 1, 1), 5); + VERIFY_IS_EQUAL(c(0, 1, 2), 6); + VERIFY_IS_EQUAL(c(0, 1, 3), 7); + VERIFY_IS_EQUAL(c(1, 0, 0), 10); + VERIFY_IS_EQUAL(c(1, 0, 1), 11); + VERIFY_IS_EQUAL(c(1, 0, 2), 12); + VERIFY_IS_EQUAL(c(1, 0, 3), 13); + VERIFY_IS_EQUAL(c(1, 1, 0), 14); + VERIFY_IS_EQUAL(c(1, 1, 1), 15); + VERIFY_IS_EQUAL(c(1, 1, 2), 16); + VERIFY_IS_EQUAL(c(1, 1, 3), 17); + VERIFY_IS_EQUAL(c(2, 0, 0), 20); + VERIFY_IS_EQUAL(c(2, 0, 1), 21); + VERIFY_IS_EQUAL(c(2, 0, 2), 22); + VERIFY_IS_EQUAL(c(2, 0, 3), 23); + VERIFY_IS_EQUAL(c(2, 1, 0), 24); + VERIFY_IS_EQUAL(c(2, 1, 1), 25); + VERIFY_IS_EQUAL(c(2, 1, 2), 26); + VERIFY_IS_EQUAL(c(2, 1, 3), 27); +#endif // EIGEN_HAS_VARIADIC_TEMPLATES +} void test_cxx11_tensor_assign() { @@ -294,4 +366,5 @@ void test_cxx11_tensor_assign() CALL_SUBTEST(test_same_type()); CALL_SUBTEST(test_auto_resize()); CALL_SUBTEST(test_compound_assign()); + CALL_SUBTEST(test_std_initializers_tensor()); } diff --git a/unsupported/test/cxx11_tensor_broadcasting.cpp b/unsupported/test/cxx11_tensor_broadcasting.cpp index 9663912a4..f0792bdcf 100644 --- a/unsupported/test/cxx11_tensor_broadcasting.cpp +++ b/unsupported/test/cxx11_tensor_broadcasting.cpp @@ -13,9 +13,10 @@ using Eigen::Tensor; +template <int DataLayout> static void test_simple_broadcasting() { - Tensor<float, 4> tensor(2,3,5,7); + Tensor<float, 4, DataLayout> tensor(2,3,5,7); tensor.setRandom(); array<ptrdiff_t, 4> broadcasts; broadcasts[0] = 1; @@ -23,7 +24,7 @@ static void test_simple_broadcasting() broadcasts[2] = 1; broadcasts[3] = 1; - Tensor<float, 4> no_broadcast; + Tensor<float, 4, DataLayout> no_broadcast; no_broadcast = tensor.broadcast(broadcasts); VERIFY_IS_EQUAL(no_broadcast.dimension(0), 2); @@ -45,7 +46,7 @@ static void test_simple_broadcasting() broadcasts[1] = 3; broadcasts[2] = 1; broadcasts[3] = 4; - Tensor<float, 4> broadcast; + Tensor<float, 4, DataLayout> broadcast; broadcast = tensor.broadcast(broadcasts); VERIFY_IS_EQUAL(broadcast.dimension(0), 4); @@ -65,16 +66,17 @@ static void test_simple_broadcasting() } +template <int DataLayout> static void test_vectorized_broadcasting() { - Tensor<float, 3> tensor(8,3,5); + Tensor<float, 3, DataLayout> tensor(8,3,5); tensor.setRandom(); array<ptrdiff_t, 3> broadcasts; broadcasts[0] = 2; broadcasts[1] = 3; broadcasts[2] = 4; - Tensor<float, 3> broadcast; + Tensor<float, 3, DataLayout> broadcast; broadcast = tensor.broadcast(broadcasts); VERIFY_IS_EQUAL(broadcast.dimension(0), 16); @@ -107,8 +109,78 @@ static void test_vectorized_broadcasting() } +template <int DataLayout> +static void test_static_broadcasting() +{ + Tensor<float, 3, DataLayout> tensor(8,3,5); + tensor.setRandom(); + Eigen::IndexList<Eigen::type2index<2>, Eigen::type2index<3>, Eigen::type2index<4>> broadcasts; + + Tensor<float, 3, DataLayout> broadcast; + broadcast = tensor.broadcast(broadcasts); + + VERIFY_IS_EQUAL(broadcast.dimension(0), 16); + VERIFY_IS_EQUAL(broadcast.dimension(1), 9); + VERIFY_IS_EQUAL(broadcast.dimension(2), 20); + + for (int i = 0; i < 16; ++i) { + for (int j = 0; j < 9; ++j) { + for (int k = 0; k < 20; ++k) { + VERIFY_IS_EQUAL(tensor(i%8,j%3,k%5), broadcast(i,j,k)); + } + } + } + + tensor.resize(11,3,5); + tensor.setRandom(); + broadcast = tensor.broadcast(broadcasts); + + VERIFY_IS_EQUAL(broadcast.dimension(0), 22); + VERIFY_IS_EQUAL(broadcast.dimension(1), 9); + VERIFY_IS_EQUAL(broadcast.dimension(2), 20); + + for (int i = 0; i < 22; ++i) { + for (int j = 0; j < 9; ++j) { + for (int k = 0; k < 20; ++k) { + VERIFY_IS_EQUAL(tensor(i%11,j%3,k%5), broadcast(i,j,k)); + } + } + } +} + + +template <int DataLayout> +static void test_fixed_size_broadcasting() +{ + // Need to add a [] operator to the Size class for this to work +#if 0 + Tensor<float, 1, DataLayout> t1(10); + t1.setRandom(); + TensorFixedSize<float, Sizes<1>, DataLayout> t2; + t2 = t2.constant(20.0f); + + Tensor<float, 1, DataLayout> t3 = t1 + t2.broadcast(Eigen::array<int, 1>{{10}}); + for (int i = 0; i < 10; ++i) { + VERIFY_IS_APPROX(t3(i), t1(i) + t2(0)); + } + + TensorMap<TensorFixedSize<float, Sizes<1>, DataLayout> > t4(t2.data(), {{1}}); + Tensor<float, 1, DataLayout> t5 = t1 + t4.broadcast(Eigen::array<int, 1>{{10}}); + for (int i = 0; i < 10; ++i) { + VERIFY_IS_APPROX(t5(i), t1(i) + t2(0)); + } +#endif +} + + void test_cxx11_tensor_broadcasting() { - CALL_SUBTEST(test_simple_broadcasting()); - CALL_SUBTEST(test_vectorized_broadcasting()); + CALL_SUBTEST(test_simple_broadcasting<ColMajor>()); + CALL_SUBTEST(test_simple_broadcasting<RowMajor>()); + CALL_SUBTEST(test_vectorized_broadcasting<ColMajor>()); + CALL_SUBTEST(test_vectorized_broadcasting<RowMajor>()); + CALL_SUBTEST(test_static_broadcasting<ColMajor>()); + CALL_SUBTEST(test_static_broadcasting<RowMajor>()); + CALL_SUBTEST(test_fixed_size_broadcasting<ColMajor>()); + CALL_SUBTEST(test_fixed_size_broadcasting<RowMajor>()); } diff --git a/unsupported/test/cxx11_tensor_chipping.cpp b/unsupported/test/cxx11_tensor_chipping.cpp index 0027b2888..0de7bbac6 100644 --- a/unsupported/test/cxx11_tensor_chipping.cpp +++ b/unsupported/test/cxx11_tensor_chipping.cpp @@ -13,18 +13,20 @@ using Eigen::Tensor; - +template<int DataLayout> static void test_simple_chip() { - Tensor<float, 5> tensor(2,3,5,7,11); + Tensor<float, 5, DataLayout> tensor(2,3,5,7,11); tensor.setRandom(); - Tensor<float, 4> chip1; - chip1 = tensor.chip<0>(1); + Tensor<float, 4, DataLayout> chip1; + chip1 = tensor.template chip<0>(1); + VERIFY_IS_EQUAL(chip1.dimension(0), 3); VERIFY_IS_EQUAL(chip1.dimension(1), 5); VERIFY_IS_EQUAL(chip1.dimension(2), 7); VERIFY_IS_EQUAL(chip1.dimension(3), 11); + for (int i = 0; i < 3; ++i) { for (int j = 0; j < 5; ++j) { for (int k = 0; k < 7; ++k) { @@ -35,7 +37,7 @@ static void test_simple_chip() } } - Tensor<float, 4> chip2 = tensor.chip<1>(1); + Tensor<float, 4, DataLayout> chip2 = tensor.template chip<1>(1); VERIFY_IS_EQUAL(chip2.dimension(0), 2); VERIFY_IS_EQUAL(chip2.dimension(1), 5); VERIFY_IS_EQUAL(chip2.dimension(2), 7); @@ -50,7 +52,7 @@ static void test_simple_chip() } } - Tensor<float, 4> chip3 = tensor.chip<2>(2); + Tensor<float, 4, DataLayout> chip3 = tensor.template chip<2>(2); VERIFY_IS_EQUAL(chip3.dimension(0), 2); VERIFY_IS_EQUAL(chip3.dimension(1), 3); VERIFY_IS_EQUAL(chip3.dimension(2), 7); @@ -65,7 +67,7 @@ static void test_simple_chip() } } - Tensor<float, 4> chip4(tensor.chip<3>(5)); + Tensor<float, 4, DataLayout> chip4(tensor.template chip<3>(5)); VERIFY_IS_EQUAL(chip4.dimension(0), 2); VERIFY_IS_EQUAL(chip4.dimension(1), 3); VERIFY_IS_EQUAL(chip4.dimension(2), 5); @@ -80,7 +82,7 @@ static void test_simple_chip() } } - Tensor<float, 4> chip5(tensor.chip<4>(7)); + Tensor<float, 4, DataLayout> chip5(tensor.template chip<4>(7)); VERIFY_IS_EQUAL(chip5.dimension(0), 2); VERIFY_IS_EQUAL(chip5.dimension(1), 3); VERIFY_IS_EQUAL(chip5.dimension(2), 5); @@ -96,14 +98,97 @@ static void test_simple_chip() } } +template<int DataLayout> +static void test_dynamic_chip() +{ + Tensor<float, 5, DataLayout> tensor(2,3,5,7,11); + tensor.setRandom(); + + Tensor<float, 4, DataLayout> chip1; + chip1 = tensor.chip(1, 0); + VERIFY_IS_EQUAL(chip1.dimension(0), 3); + VERIFY_IS_EQUAL(chip1.dimension(1), 5); + VERIFY_IS_EQUAL(chip1.dimension(2), 7); + VERIFY_IS_EQUAL(chip1.dimension(3), 11); + for (int i = 0; i < 3; ++i) { + for (int j = 0; j < 5; ++j) { + for (int k = 0; k < 7; ++k) { + for (int l = 0; l < 11; ++l) { + VERIFY_IS_EQUAL(chip1(i,j,k,l), tensor(1,i,j,k,l)); + } + } + } + } + + Tensor<float, 4, DataLayout> chip2 = tensor.chip(1, 1); + VERIFY_IS_EQUAL(chip2.dimension(0), 2); + VERIFY_IS_EQUAL(chip2.dimension(1), 5); + VERIFY_IS_EQUAL(chip2.dimension(2), 7); + VERIFY_IS_EQUAL(chip2.dimension(3), 11); + for (int i = 0; i < 2; ++i) { + for (int j = 0; j < 3; ++j) { + for (int k = 0; k < 7; ++k) { + for (int l = 0; l < 11; ++l) { + VERIFY_IS_EQUAL(chip2(i,j,k,l), tensor(i,1,j,k,l)); + } + } + } + } + + Tensor<float, 4, DataLayout> chip3 = tensor.chip(2, 2); + VERIFY_IS_EQUAL(chip3.dimension(0), 2); + VERIFY_IS_EQUAL(chip3.dimension(1), 3); + VERIFY_IS_EQUAL(chip3.dimension(2), 7); + VERIFY_IS_EQUAL(chip3.dimension(3), 11); + for (int i = 0; i < 2; ++i) { + for (int j = 0; j < 3; ++j) { + for (int k = 0; k < 7; ++k) { + for (int l = 0; l < 11; ++l) { + VERIFY_IS_EQUAL(chip3(i,j,k,l), tensor(i,j,2,k,l)); + } + } + } + } + + Tensor<float, 4, DataLayout> chip4(tensor.chip(5, 3)); + VERIFY_IS_EQUAL(chip4.dimension(0), 2); + VERIFY_IS_EQUAL(chip4.dimension(1), 3); + VERIFY_IS_EQUAL(chip4.dimension(2), 5); + VERIFY_IS_EQUAL(chip4.dimension(3), 11); + for (int i = 0; i < 2; ++i) { + for (int j = 0; j < 3; ++j) { + for (int k = 0; k < 5; ++k) { + for (int l = 0; l < 7; ++l) { + VERIFY_IS_EQUAL(chip4(i,j,k,l), tensor(i,j,k,5,l)); + } + } + } + } + + Tensor<float, 4, DataLayout> chip5(tensor.chip(7, 4)); + VERIFY_IS_EQUAL(chip5.dimension(0), 2); + VERIFY_IS_EQUAL(chip5.dimension(1), 3); + VERIFY_IS_EQUAL(chip5.dimension(2), 5); + VERIFY_IS_EQUAL(chip5.dimension(3), 7); + for (int i = 0; i < 2; ++i) { + for (int j = 0; j < 3; ++j) { + for (int k = 0; k < 5; ++k) { + for (int l = 0; l < 7; ++l) { + VERIFY_IS_EQUAL(chip5(i,j,k,l), tensor(i,j,k,l,7)); + } + } + } + } +} +template<int DataLayout> static void test_chip_in_expr() { - Tensor<float, 5> input1(2,3,5,7,11); + Tensor<float, 5, DataLayout> input1(2,3,5,7,11); input1.setRandom(); - Tensor<float, 4> input2(3,5,7,11); + Tensor<float, 4, DataLayout> input2(3,5,7,11); input2.setRandom(); - Tensor<float, 4> result = input1.chip<0>(0) + input2; + Tensor<float, 4, DataLayout> result = input1.template chip<0>(0) + input2; for (int i = 0; i < 3; ++i) { for (int j = 0; j < 5; ++j) { for (int k = 0; k < 7; ++k) { @@ -115,9 +200,9 @@ static void test_chip_in_expr() { } } - Tensor<float, 3> input3(3,7,11); + Tensor<float, 3, DataLayout> input3(3,7,11); input3.setRandom(); - Tensor<float, 3> result2 = input1.chip<0>(0).chip<1>(2) + input3; + Tensor<float, 3, DataLayout> result2 = input1.template chip<0>(0).template chip<1>(2) + input3; for (int i = 0; i < 3; ++i) { for (int j = 0; j < 7; ++j) { for (int k = 0; k < 11; ++k) { @@ -128,16 +213,16 @@ static void test_chip_in_expr() { } } - +template<int DataLayout> static void test_chip_as_lvalue() { - Tensor<float, 5> input1(2,3,5,7,11); + Tensor<float, 5, DataLayout> input1(2,3,5,7,11); input1.setRandom(); - Tensor<float, 4> input2(3,5,7,11); + Tensor<float, 4, DataLayout> input2(3,5,7,11); input2.setRandom(); - Tensor<float, 5> tensor = input1; - tensor.chip<0>(1) = input2; + Tensor<float, 5, DataLayout> tensor = input1; + tensor.template chip<0>(1) = input2; for (int i = 0; i < 2; ++i) { for (int j = 0; j < 3; ++j) { for (int k = 0; k < 5; ++k) { @@ -154,10 +239,10 @@ static void test_chip_as_lvalue() } } - Tensor<float, 4> input3(2,5,7,11); + Tensor<float, 4, DataLayout> input3(2,5,7,11); input3.setRandom(); tensor = input1; - tensor.chip<1>(1) = input3; + tensor.template chip<1>(1) = input3; for (int i = 0; i < 2; ++i) { for (int j = 0; j < 3; ++j) { for (int k = 0; k < 5; ++k) { @@ -174,10 +259,10 @@ static void test_chip_as_lvalue() } } - Tensor<float, 4> input4(2,3,7,11); + Tensor<float, 4, DataLayout> input4(2,3,7,11); input4.setRandom(); tensor = input1; - tensor.chip<2>(3) = input4; + tensor.template chip<2>(3) = input4; for (int i = 0; i < 2; ++i) { for (int j = 0; j < 3; ++j) { for (int k = 0; k < 5; ++k) { @@ -194,10 +279,10 @@ static void test_chip_as_lvalue() } } - Tensor<float, 4> input5(2,3,5,11); + Tensor<float, 4, DataLayout> input5(2,3,5,11); input5.setRandom(); tensor = input1; - tensor.chip<3>(4) = input5; + tensor.template chip<3>(4) = input5; for (int i = 0; i < 2; ++i) { for (int j = 0; j < 3; ++j) { for (int k = 0; k < 5; ++k) { @@ -214,10 +299,10 @@ static void test_chip_as_lvalue() } } - Tensor<float, 4> input6(2,3,5,7); + Tensor<float, 4, DataLayout> input6(2,3,5,7); input6.setRandom(); tensor = input1; - tensor.chip<4>(5) = input6; + tensor.template chip<4>(5) = input6; for (int i = 0; i < 2; ++i) { for (int j = 0; j < 3; ++j) { for (int k = 0; k < 5; ++k) { @@ -235,47 +320,57 @@ static void test_chip_as_lvalue() } } - +template<int DataLayout> static void test_chip_raw_data() { - Tensor<float, 5> tensor(2,3,5,7,11); + Tensor<float, 5, DataLayout> tensor(2,3,5,7,11); tensor.setRandom(); - typedef TensorEvaluator<decltype(tensor.chip<4>(3)), DefaultDevice> Evaluator4; - auto chip = Evaluator4(tensor.chip<4>(3), DefaultDevice()); + typedef TensorEvaluator<decltype(tensor.template chip<4>(3)), DefaultDevice> Evaluator4; + auto chip = Evaluator4(tensor.template chip<4>(3), DefaultDevice()); for (int i = 0; i < 2; ++i) { for (int j = 0; j < 3; ++j) { for (int k = 0; k < 5; ++k) { for (int l = 0; l < 7; ++l) { - int chip_index = i + 2 * (j + 3 * (k + 5 * l)); + int chip_index; + if (DataLayout == ColMajor) { + chip_index = i + 2 * (j + 3 * (k + 5 * l)); + } else { + chip_index = 11 * (l + 7 * (k + 5 * (j + 3 * i))); + } VERIFY_IS_EQUAL(chip.data()[chip_index], tensor(i,j,k,l,3)); } } } } - typedef TensorEvaluator<decltype(tensor.chip<0>(0)), DefaultDevice> Evaluator0; - auto chip0 = Evaluator0(tensor.chip<0>(0), DefaultDevice()); + typedef TensorEvaluator<decltype(tensor.template chip<0>(0)), DefaultDevice> Evaluator0; + auto chip0 = Evaluator0(tensor.template chip<0>(0), DefaultDevice()); VERIFY_IS_EQUAL(chip0.data(), static_cast<float*>(0)); - typedef TensorEvaluator<decltype(tensor.chip<1>(0)), DefaultDevice> Evaluator1; - auto chip1 = Evaluator1(tensor.chip<1>(0), DefaultDevice()); + typedef TensorEvaluator<decltype(tensor.template chip<1>(0)), DefaultDevice> Evaluator1; + auto chip1 = Evaluator1(tensor.template chip<1>(0), DefaultDevice()); VERIFY_IS_EQUAL(chip1.data(), static_cast<float*>(0)); - typedef TensorEvaluator<decltype(tensor.chip<2>(0)), DefaultDevice> Evaluator2; - auto chip2 = Evaluator2(tensor.chip<2>(0), DefaultDevice()); + typedef TensorEvaluator<decltype(tensor.template chip<2>(0)), DefaultDevice> Evaluator2; + auto chip2 = Evaluator2(tensor.template chip<2>(0), DefaultDevice()); VERIFY_IS_EQUAL(chip2.data(), static_cast<float*>(0)); - typedef TensorEvaluator<decltype(tensor.chip<3>(0)), DefaultDevice> Evaluator3; - auto chip3 = Evaluator3(tensor.chip<3>(0), DefaultDevice()); + typedef TensorEvaluator<decltype(tensor.template chip<3>(0)), DefaultDevice> Evaluator3; + auto chip3 = Evaluator3(tensor.template chip<3>(0), DefaultDevice()); VERIFY_IS_EQUAL(chip3.data(), static_cast<float*>(0)); } - void test_cxx11_tensor_chipping() { - CALL_SUBTEST(test_simple_chip()); - CALL_SUBTEST(test_chip_in_expr()); - CALL_SUBTEST(test_chip_as_lvalue()); - CALL_SUBTEST(test_chip_raw_data()); + CALL_SUBTEST(test_simple_chip<ColMajor>()); + CALL_SUBTEST(test_simple_chip<RowMajor>()); + CALL_SUBTEST(test_dynamic_chip<ColMajor>()); + CALL_SUBTEST(test_dynamic_chip<RowMajor>()); + CALL_SUBTEST(test_chip_in_expr<ColMajor>()); + CALL_SUBTEST(test_chip_in_expr<RowMajor>()); + CALL_SUBTEST(test_chip_as_lvalue<ColMajor>()); + CALL_SUBTEST(test_chip_as_lvalue<RowMajor>()); + CALL_SUBTEST(test_chip_raw_data<ColMajor>()); + CALL_SUBTEST(test_chip_raw_data<RowMajor>()); } diff --git a/unsupported/test/cxx11_tensor_concatenation.cpp b/unsupported/test/cxx11_tensor_concatenation.cpp index 8fd4f5f80..9fdf33c16 100644 --- a/unsupported/test/cxx11_tensor_concatenation.cpp +++ b/unsupported/test/cxx11_tensor_concatenation.cpp @@ -13,15 +13,16 @@ using Eigen::Tensor; +template<int DataLayout> static void test_dimension_failures() { - Tensor<int, 3> left(2, 3, 1); - Tensor<int, 3> right(3, 3, 1); + Tensor<int, 3, DataLayout> left(2, 3, 1); + Tensor<int, 3, DataLayout> right(3, 3, 1); left.setRandom(); right.setRandom(); // Okay; other dimensions are equal. - Tensor<int, 3> concatenation = left.concatenate(right, 0); + Tensor<int, 3, DataLayout> concatenation = left.concatenate(right, 0); // Dimension mismatches. VERIFY_RAISES_ASSERT(concatenation = left.concatenate(right, 1)); @@ -32,33 +33,35 @@ static void test_dimension_failures() VERIFY_RAISES_ASSERT(concatenation = left.concatenate(right, -1)); } +template<int DataLayout> static void test_static_dimension_failure() { - Tensor<int, 2> left(2, 3); - Tensor<int, 3> right(2, 3, 1); + Tensor<int, 2, DataLayout> left(2, 3); + Tensor<int, 3, DataLayout> right(2, 3, 1); #ifdef CXX11_TENSOR_CONCATENATION_STATIC_DIMENSION_FAILURE // Technically compatible, but we static assert that the inputs have same // NumDims. - Tensor<int, 3> concatenation = left.concatenate(right, 0); + Tensor<int, 3, DataLayout> concatenation = left.concatenate(right, 0); #endif // This can be worked around in this case. - Tensor<int, 3> concatenation = left + Tensor<int, 3, DataLayout> concatenation = left .reshape(Tensor<int, 3>::Dimensions{{2, 3, 1}}) .concatenate(right, 0); - Tensor<int, 2> alternative = left + Tensor<int, 2, DataLayout> alternative = left .concatenate(right.reshape(Tensor<int, 2>::Dimensions{{2, 3}}), 0); } +template<int DataLayout> static void test_simple_concatenation() { - Tensor<int, 3> left(2, 3, 1); - Tensor<int, 3> right(2, 3, 1); + Tensor<int, 3, DataLayout> left(2, 3, 1); + Tensor<int, 3, DataLayout> right(2, 3, 1); left.setRandom(); right.setRandom(); - Tensor<int, 3> concatenation = left.concatenate(right, 0); + Tensor<int, 3, DataLayout> concatenation = left.concatenate(right, 0); VERIFY_IS_EQUAL(concatenation.dimension(0), 4); VERIFY_IS_EQUAL(concatenation.dimension(1), 3); VERIFY_IS_EQUAL(concatenation.dimension(2), 1); @@ -103,8 +106,11 @@ static void test_simple_concatenation() void test_cxx11_tensor_concatenation() { - CALL_SUBTEST(test_dimension_failures()); - CALL_SUBTEST(test_static_dimension_failure()); - CALL_SUBTEST(test_simple_concatenation()); + CALL_SUBTEST(test_dimension_failures<ColMajor>()); + CALL_SUBTEST(test_dimension_failures<RowMajor>()); + CALL_SUBTEST(test_static_dimension_failure<ColMajor>()); + CALL_SUBTEST(test_static_dimension_failure<RowMajor>()); + CALL_SUBTEST(test_simple_concatenation<ColMajor>()); + CALL_SUBTEST(test_simple_concatenation<RowMajor>()); // CALL_SUBTEST(test_vectorized_concatenation()); } diff --git a/unsupported/test/cxx11_tensor_contract_cuda.cpp b/unsupported/test/cxx11_tensor_contract_cuda.cpp new file mode 100644 index 000000000..9599607c6 --- /dev/null +++ b/unsupported/test/cxx11_tensor_contract_cuda.cpp @@ -0,0 +1,121 @@ +// This file is part of Eigen, a lightweight C++ template library +// for linear algebra. +// +// Copyright (C) 2014 Benoit Steiner <benoit.steiner.goog@gmail.com> +// Copyright (C) 2014 Navdeep Jaitly <ndjaitly@google.com> +// +// 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/. + +#define EIGEN_TEST_NO_LONGDOUBLE +#define EIGEN_TEST_NO_COMPLEX +#define EIGEN_TEST_FUNC cxx11_tensor_cuda +#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int +#define EIGEN_USE_GPU + + +#include "main.h" +#include <unsupported/Eigen/CXX11/Tensor> + +using Eigen::Tensor; +typedef Tensor<float, 1>::DimensionPair DimPair; + +template<int DataLayout> +static void test_cuda_contraction(int m_size, int k_size, int n_size) +{ + cout<<"Calling with ("<<m_size<<","<<k_size<<","<<n_size<<")"<<std::endl; + // with these dimensions, the output has 300 * 140 elements, which is + // more than 30 * 1024, which is the number of threads in blocks on + // a 15 SM GK110 GPU + Tensor<float, 2, DataLayout> t_left(Eigen::array<int, 2>(m_size, k_size)); + Tensor<float, 2, DataLayout> t_right(Eigen::array<int, 2>(k_size, n_size)); + Tensor<float, 2, DataLayout> t_result(Eigen::array<int, 2>(m_size, n_size)); + Tensor<float, 2, DataLayout> t_result_gpu(Eigen::array<int, 2>(m_size, n_size)); + Eigen::array<DimPair, 1> dims(DimPair(1, 0)); + + t_left.setRandom(); + t_right.setRandom(); + + std::size_t t_left_bytes = t_left.size() * sizeof(float); + std::size_t t_right_bytes = t_right.size() * sizeof(float); + std::size_t t_result_bytes = t_result.size() * sizeof(float); + + float* d_t_left; + float* d_t_right; + float* d_t_result; + + cudaMalloc((void**)(&d_t_left), t_left_bytes); + cudaMalloc((void**)(&d_t_right), t_right_bytes); + cudaMalloc((void**)(&d_t_result), t_result_bytes); + + cudaMemcpy(d_t_left, t_left.data(), t_left_bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_t_right, t_right.data(), t_right_bytes, cudaMemcpyHostToDevice); + + cudaStream_t stream; + assert(cudaStreamCreate(&stream) == cudaSuccess); + Eigen::GpuDevice gpu_device(&stream); + + Eigen::TensorMap<Eigen::Tensor<float, 2, DataLayout> > + gpu_t_left(d_t_left, Eigen::array<int, 2>(m_size, k_size)); + Eigen::TensorMap<Eigen::Tensor<float, 2, DataLayout> > + gpu_t_right(d_t_right, Eigen::array<int, 2>(k_size, n_size)); + Eigen::TensorMap<Eigen::Tensor<float, 2, DataLayout> > + gpu_t_result(d_t_result, Eigen::array<int, 2>(m_size, n_size)); + + + gpu_t_result.device(gpu_device) = gpu_t_left.contract(gpu_t_right, dims); + t_result = t_left.contract(t_right, dims); + + cudaMemcpy(t_result_gpu.data(), d_t_result, t_result_bytes, cudaMemcpyDeviceToHost); + for (size_t i = 0; i < t_result.dimensions().TotalSize(); i++) { + if (fabs(t_result.data()[i] - t_result_gpu.data()[i]) >= 1e-4) { + cout << "mismatch detected at index " << i << ": " << t_result.data()[i] + << " vs " << t_result_gpu.data()[i] << endl; + assert(false); + } + } + + cudaFree((void*)d_t_left); + cudaFree((void*)d_t_right); + cudaFree((void*)d_t_result); +} + + +void test_cxx11_tensor_cuda() +{ + cout<<"Calling contraction tests"<<std::endl; + CALL_SUBTEST(test_cuda_contraction<ColMajor>(128, 128, 128)); + CALL_SUBTEST(test_cuda_contraction<RowMajor>(128, 128, 128)); + for (int k = 32; k < 256; k++) { + CALL_SUBTEST(test_cuda_contraction<ColMajor>(128, k, 128)); + CALL_SUBTEST(test_cuda_contraction<RowMajor>(128, k, 128)); + } + for (int k = 32; k < 256; k++) { + CALL_SUBTEST(test_cuda_contraction<ColMajor>(128, 128, k)); + CALL_SUBTEST(test_cuda_contraction<RowMajor>(128, 128, k)); + } + for (int k = 32; k < 256; k++) { + CALL_SUBTEST(test_cuda_contraction<ColMajor>(k, 128, 128)); + CALL_SUBTEST(test_cuda_contraction<RowMajor>(k, 128, 128)); + } + + int m_sizes[] = {31, 39, 63, 64, 65, + 127, 129, 255, 257, 511, + 512, 513, 1023, 1024, 1025 }; + int n_sizes[] = {31, 39, 63, 64, 65, + 127, 129, 255, 257, 511, + 512, 513, 1023, 1024, 1025 }; + + int k_sizes[] = { 31, 39, 63, 64, 65, + 95, 96, 127, 129, 255, + 257, 511, 512, 513, 1023, + 1024, 1025}; + + for (int i = 0; i <15; i++) + for (int j = 0; j < 15; j++) + for (int k = 0; k < 17; k++) { + CALL_SUBTEST(test_cuda_contraction<ColMajor>(m_sizes[i], n_sizes[j], k_sizes[k])); + CALL_SUBTEST(test_cuda_contraction<RowMajor>(m_sizes[i], n_sizes[j], k_sizes[k])); + } +} diff --git a/unsupported/test/cxx11_tensor_contraction.cpp b/unsupported/test/cxx11_tensor_contraction.cpp index 17bd335f7..6124818fd 100644 --- a/unsupported/test/cxx11_tensor_contraction.cpp +++ b/unsupported/test/cxx11_tensor_contraction.cpp @@ -16,18 +16,18 @@ using Eigen::Tensor; typedef Tensor<float, 1>::DimensionPair DimPair; - +template<int DataLayout> static void test_evals() { - Tensor<float, 2> mat1(2, 3); - Tensor<float, 2> mat2(2, 3); - Tensor<float, 2> mat3(3, 2); + Tensor<float, 2, DataLayout> mat1(2, 3); + Tensor<float, 2, DataLayout> mat2(2, 3); + Tensor<float, 2, DataLayout> mat3(3, 2); mat1.setRandom(); mat2.setRandom(); mat3.setRandom(); - Tensor<float, 2> mat4(3,3); + Tensor<float, 2, DataLayout> mat4(3,3); mat4.setZero(); Eigen::array<DimPair, 1> dims3({{DimPair(0, 0)}}); typedef TensorEvaluator<decltype(mat1.contract(mat2, dims3)), DefaultDevice> Evaluator; @@ -47,7 +47,7 @@ static void test_evals() VERIFY_IS_APPROX(mat4(2,1), mat1(0,2)*mat2(0,1) + mat1(1,2)*mat2(1,1)); VERIFY_IS_APPROX(mat4(2,2), mat1(0,2)*mat2(0,2) + mat1(1,2)*mat2(1,2)); - Tensor<float, 2> mat5(2,2); + Tensor<float, 2, DataLayout> mat5(2,2); mat5.setZero(); Eigen::array<DimPair, 1> dims4({{DimPair(1, 1)}}); typedef TensorEvaluator<decltype(mat1.contract(mat2, dims4)), DefaultDevice> Evaluator2; @@ -62,7 +62,7 @@ static void test_evals() VERIFY_IS_APPROX(mat5(1,0), mat1(1,0)*mat2(0,0) + mat1(1,1)*mat2(0,1) + mat1(1,2)*mat2(0,2)); VERIFY_IS_APPROX(mat5(1,1), mat1(1,0)*mat2(1,0) + mat1(1,1)*mat2(1,1) + mat1(1,2)*mat2(1,2)); - Tensor<float, 2> mat6(2,2); + Tensor<float, 2, DataLayout> mat6(2,2); mat6.setZero(); Eigen::array<DimPair, 1> dims6({{DimPair(1, 0)}}); typedef TensorEvaluator<decltype(mat1.contract(mat3, dims6)), DefaultDevice> Evaluator3; @@ -78,16 +78,16 @@ static void test_evals() VERIFY_IS_APPROX(mat6(1,1), mat1(1,0)*mat3(0,1) + mat1(1,1)*mat3(1,1) + mat1(1,2)*mat3(2,1)); } - +template<int DataLayout> static void test_scalar() { - Tensor<float, 1> vec1({6}); - Tensor<float, 1> vec2({6}); + Tensor<float, 1, DataLayout> vec1({6}); + Tensor<float, 1, DataLayout> vec2({6}); vec1.setRandom(); vec2.setRandom(); - Tensor<float, 1> scalar(1); + Tensor<float, 1, DataLayout> scalar(1); scalar.setZero(); Eigen::array<DimPair, 1> dims({{DimPair(0, 0)}}); typedef TensorEvaluator<decltype(vec1.contract(vec2, dims)), DefaultDevice> Evaluator; @@ -102,16 +102,16 @@ static void test_scalar() VERIFY_IS_APPROX(scalar(0), expected); } - +template<int DataLayout> static void test_multidims() { - Tensor<float, 3> mat1(2, 2, 2); - Tensor<float, 4> mat2(2, 2, 2, 2); + Tensor<float, 3, DataLayout> mat1(2, 2, 2); + Tensor<float, 4, DataLayout> mat2(2, 2, 2, 2); mat1.setRandom(); mat2.setRandom(); - Tensor<float, 3> mat3(2, 2, 2); + Tensor<float, 3, DataLayout> mat3(2, 2, 2); mat3.setZero(); Eigen::array<DimPair, 2> dims({{DimPair(1, 2), DimPair(2, 3)}}); typedef TensorEvaluator<decltype(mat1.contract(mat2, dims)), DefaultDevice> Evaluator; @@ -140,15 +140,15 @@ static void test_multidims() mat1(1,0,1)*mat2(1,1,0,1) + mat1(1,1,1)*mat2(1,1,1,1)); } - +template<int DataLayout> static void test_holes() { - Tensor<float, 4> t1(2, 5, 7, 3); - Tensor<float, 5> t2(2, 7, 11, 13, 3); + Tensor<float, 4, DataLayout> t1(2, 5, 7, 3); + Tensor<float, 5, DataLayout> t2(2, 7, 11, 13, 3); t1.setRandom(); t2.setRandom(); Eigen::array<DimPair, 2> dims({{DimPair(0, 0), DimPair(3, 4)}}); - Tensor<float, 5> result = t1.contract(t2, dims); + Tensor<float, 5, DataLayout> result = t1.contract(t2, dims); VERIFY_IS_EQUAL(result.dimension(0), 5); VERIFY_IS_EQUAL(result.dimension(1), 7); VERIFY_IS_EQUAL(result.dimension(2), 7); @@ -174,16 +174,16 @@ static void test_holes() { } } - +template<int DataLayout> static void test_full_redux() { - Tensor<float, 2> t1(2, 2); - Tensor<float, 3> t2(2, 2, 2); + Tensor<float, 2, DataLayout> t1(2, 2); + Tensor<float, 3, DataLayout> t2(2, 2, 2); t1.setRandom(); t2.setRandom(); Eigen::array<DimPair, 2> dims({{DimPair(0, 0), DimPair(1, 1)}}); - Tensor<float, 1> result = t1.contract(t2, dims); + Tensor<float, 1, DataLayout> result = t1.contract(t2, dims); VERIFY_IS_EQUAL(result.dimension(0), 2); VERIFY_IS_APPROX(result(0), t1(0, 0) * t2(0, 0, 0) + t1(1, 0) * t2(1, 0, 0) + t1(0, 1) * t2(0, 1, 0) + t1(1, 1) * t2(1, 1, 0)); @@ -200,13 +200,13 @@ static void test_full_redux() + t1(0, 1) * t2(1, 0, 1) + t1(1, 1) * t2(1, 1, 1)); } - +template<int DataLayout> static void test_contraction_of_contraction() { - Tensor<float, 2> t1(2, 2); - Tensor<float, 2> t2(2, 2); - Tensor<float, 2> t3(2, 2); - Tensor<float, 2> t4(2, 2); + Tensor<float, 2, DataLayout> t1(2, 2); + Tensor<float, 2, DataLayout> t2(2, 2); + Tensor<float, 2, DataLayout> t3(2, 2); + Tensor<float, 2, DataLayout> t4(2, 2); t1.setRandom(); t2.setRandom(); t3.setRandom(); @@ -216,30 +216,32 @@ static void test_contraction_of_contraction() auto contract1 = t1.contract(t2, dims); auto diff = t3 - contract1; auto contract2 = t1.contract(t4, dims); - Tensor<float, 2> result = contract2.contract(diff, dims); + Tensor<float, 2, DataLayout> result = contract2.contract(diff, dims); + VERIFY_IS_EQUAL(result.dimension(0), 2); VERIFY_IS_EQUAL(result.dimension(1), 2); - Eigen::Map<MatrixXf> m1(t1.data(), 2, 2); - Eigen::Map<MatrixXf> m2(t2.data(), 2, 2); - Eigen::Map<MatrixXf> m3(t3.data(), 2, 2); - Eigen::Map<MatrixXf> m4(t4.data(), 2, 2); - Eigen::MatrixXf expected = (m1 * m4) * (m3 - m1 * m2); + Eigen::Map<Eigen::Matrix<float, Dynamic, Dynamic, DataLayout>> + m1(t1.data(), 2, 2), m2(t2.data(), 2, 2), m3(t3.data(), 2, 2), + m4(t4.data(), 2, 2); + Eigen::Matrix<float, Dynamic, Dynamic, DataLayout> + expected = (m1 * m4) * (m3 - m1 * m2); + VERIFY_IS_APPROX(result(0, 0), expected(0, 0)); VERIFY_IS_APPROX(result(0, 1), expected(0, 1)); VERIFY_IS_APPROX(result(1, 0), expected(1, 0)); VERIFY_IS_APPROX(result(1, 1), expected(1, 1)); } - +template<int DataLayout> static void test_expr() { - Tensor<float, 2> mat1(2, 3); - Tensor<float, 2> mat2(3, 2); + Tensor<float, 2, DataLayout> mat1(2, 3); + Tensor<float, 2, DataLayout> mat2(3, 2); mat1.setRandom(); mat2.setRandom(); - Tensor<float, 2> mat3(2,2); + Tensor<float, 2, DataLayout> mat3(2,2); Eigen::array<DimPair, 1> dims({{DimPair(1, 0)}}); mat3 = mat1.contract(mat2, dims); @@ -250,16 +252,16 @@ static void test_expr() VERIFY_IS_APPROX(mat3(1,1), mat1(1,0)*mat2(0,1) + mat1(1,1)*mat2(1,1) + mat1(1,2)*mat2(2,1)); } - +template<int DataLayout> static void test_out_of_order_contraction() { - Tensor<float, 3> mat1(2, 2, 2); - Tensor<float, 3> mat2(2, 2, 2); + Tensor<float, 3, DataLayout> mat1(2, 2, 2); + Tensor<float, 3, DataLayout> mat2(2, 2, 2); mat1.setRandom(); mat2.setRandom(); - Tensor<float, 2> mat3(2, 2); + Tensor<float, 2, DataLayout> mat3(2, 2); Eigen::array<DimPair, 2> dims({{DimPair(2, 0), DimPair(0, 2)}}); mat3 = mat1.contract(mat2, dims); @@ -295,18 +297,18 @@ static void test_out_of_order_contraction() } - +template<int DataLayout> static void test_consistency() { // this does something like testing (A*B)^T = (B^T * A^T) - Tensor<float, 3> mat1(4, 3, 5); - Tensor<float, 5> mat2(3, 2, 1, 5, 4); + Tensor<float, 3, DataLayout> mat1(4, 3, 5); + Tensor<float, 5, DataLayout> mat2(3, 2, 1, 5, 4); mat1.setRandom(); mat2.setRandom(); - Tensor<float, 4> mat3(5, 2, 1, 5); - Tensor<float, 4> mat4(2, 1, 5, 5); + Tensor<float, 4, DataLayout> mat3(5, 2, 1, 5); + Tensor<float, 4, DataLayout> mat4(2, 1, 5, 5); // contract on dimensions of size 4 and 3 Eigen::array<DimPair, 2> dims1({{DimPair(0, 4), DimPair(1, 0)}}); @@ -316,27 +318,40 @@ static void test_consistency() mat4 = mat2.contract(mat1, dims2); // check that these are equal except for ordering of dimensions - for (size_t i = 0; i < 5; i++) { - for (size_t j = 0; j < 10; j++) { - VERIFY_IS_APPROX(mat3.data()[i + 5 * j], mat4.data()[j + 10 * i]); + if (DataLayout == ColMajor) { + for (size_t i = 0; i < 5; i++) { + for (size_t j = 0; j < 10; j++) { + VERIFY_IS_APPROX(mat3.data()[i + 5 * j], mat4.data()[j + 10 * i]); + } + } + } else { + // Row major + for (size_t i = 0; i < 5; i++) { + for (size_t j = 0; j < 10; j++) { + VERIFY_IS_APPROX(mat3.data()[10 * i + j], mat4.data()[i + 5 * j]); + } } } } - +template<int DataLayout> static void test_large_contraction() { - Tensor<float, 4> t_left(30, 50, 8, 31); - Tensor<float, 5> t_right(8, 31, 7, 20, 10); - Tensor<float, 5> t_result(30, 50, 7, 20, 10); + Tensor<float, 4, DataLayout> t_left(30, 50, 8, 31); + Tensor<float, 5, DataLayout> t_right(8, 31, 7, 20, 10); + Tensor<float, 5, DataLayout> t_result(30, 50, 7, 20, 10); t_left.setRandom(); t_right.setRandom(); - typedef Map<MatrixXf> MapXf; + // Add a little offset so that the results won't be close to zero. + t_left += t_left.constant(1.0f); + t_right += t_right.constant(1.0f); + + typedef Map<Eigen::Matrix<float, Dynamic, Dynamic, DataLayout>> MapXf; MapXf m_left(t_left.data(), 1500, 248); MapXf m_right(t_right.data(), 248, 1400); - MatrixXf m_result(1500, 1400); + Eigen::Matrix<float, Dynamic, Dynamic, DataLayout> m_result(1500, 1400); // this contraction should be equivalent to a single matrix multiplication Eigen::array<DimPair, 2> dims({{DimPair(2, 0), DimPair(3, 1)}}); @@ -351,20 +366,20 @@ static void test_large_contraction() } } - +template<int DataLayout> static void test_matrix_vector() { - Tensor<float, 2> t_left(30, 50); - Tensor<float, 1> t_right(50); - Tensor<float, 1> t_result(30); + Tensor<float, 2, DataLayout> t_left(30, 50); + Tensor<float, 1, DataLayout> t_right(50); + Tensor<float, 1, DataLayout> t_result(30); t_left.setRandom(); t_right.setRandom(); - typedef Map<Eigen::Matrix<float, Dynamic, Dynamic>> MapXf; + typedef Map<Eigen::Matrix<float, Dynamic, Dynamic, DataLayout>> MapXf; MapXf m_left(t_left.data(), 30, 50); MapXf m_right(t_right.data(), 50, 1); - Eigen::Matrix<float, Dynamic, Dynamic> m_result(30, 1); + Eigen::Matrix<float, Dynamic, Dynamic, DataLayout> m_result(30, 1); // this contraction should be equivalent to a single matrix multiplication Eigen::array<DimPair, 1> dims{{DimPair(1, 0)}}; @@ -379,18 +394,19 @@ static void test_matrix_vector() } +template<int DataLayout> static void test_tensor_vector() { - Tensor<float, 3> t_left(7, 13, 17); - Tensor<float, 2> t_right(1, 7); - typedef typename Tensor<float, 1>::DimensionPair DimensionPair; + Tensor<float, 3, DataLayout> t_left(7, 13, 17); + Tensor<float, 2, DataLayout> t_right(1, 7); + typedef typename Tensor<float, 1, DataLayout>::DimensionPair DimensionPair; Eigen::array<DimensionPair, 1> dim_pair01{{{0, 1}}}; - Tensor<float, 3> t_result = t_left.contract(t_right, dim_pair01); + Tensor<float, 3, DataLayout> t_result = t_left.contract(t_right, dim_pair01); - typedef Map<Eigen::Matrix<float, Dynamic, Dynamic>> MapXf; + typedef Map<Eigen::Matrix<float, Dynamic, Dynamic, DataLayout>> MapXf; MapXf m_left(t_left.data(), 7, 13*17); MapXf m_right(t_right.data(), 1, 7); - Eigen::Matrix<float, Dynamic, Dynamic> m_result = m_left.transpose() * m_right.transpose(); + Eigen::Matrix<float, Dynamic, Dynamic, DataLayout> m_result = m_left.transpose() * m_right.transpose(); for (size_t i = 0; i < t_result.dimensions().TotalSize(); i++) { VERIFY_IS_APPROX(t_result(i), m_result(i, 0)); @@ -398,18 +414,63 @@ static void test_tensor_vector() } +template<int DataLayout> +static void test_small_blocking_factors() +{ + Tensor<float, 4, DataLayout> t_left(30, 5, 3, 31); + Tensor<float, 5, DataLayout> t_right(3, 31, 7, 20, 1); + t_left.setRandom(); + t_right.setRandom(); + + // Add a little offset so that the results won't be close to zero. + t_left += t_left.constant(1.0f); + t_right += t_right.constant(1.0f); + + // Force the cache sizes, which results in smaller blocking factors. + Eigen::setCpuCacheSizes(896, 1920, 2944); + + // this contraction should be equivalent to a single matrix multiplication + Eigen::array<DimPair, 2> dims({{DimPair(2, 0), DimPair(3, 1)}}); + Tensor<float, 5, DataLayout> t_result; + t_result = t_left.contract(t_right, dims); + + // compute result using a simple eigen matrix product + Map<Eigen::Matrix<float, Dynamic, Dynamic, DataLayout>> m_left(t_left.data(), 150, 93); + Map<Eigen::Matrix<float, Dynamic, Dynamic, DataLayout>> m_right(t_right.data(), 93, 140); + Eigen::Matrix<float, Dynamic, Dynamic, DataLayout> m_result = m_left * m_right; + + for (size_t i = 0; i < t_result.dimensions().TotalSize(); i++) { + VERIFY_IS_APPROX(t_result.data()[i], m_result.data()[i]); + } +} + + void test_cxx11_tensor_contraction() { - CALL_SUBTEST(test_evals()); - CALL_SUBTEST(test_scalar()); - CALL_SUBTEST(test_multidims()); - CALL_SUBTEST(test_holes()); - CALL_SUBTEST(test_full_redux()); - CALL_SUBTEST(test_contraction_of_contraction()); - CALL_SUBTEST(test_expr()); - CALL_SUBTEST(test_out_of_order_contraction()); - CALL_SUBTEST(test_consistency()); - CALL_SUBTEST(test_large_contraction()); - CALL_SUBTEST(test_matrix_vector()); - CALL_SUBTEST(test_tensor_vector()); + CALL_SUBTEST(test_evals<ColMajor>()); + CALL_SUBTEST(test_evals<RowMajor>()); + CALL_SUBTEST(test_scalar<ColMajor>()); + CALL_SUBTEST(test_scalar<RowMajor>()); + CALL_SUBTEST(test_multidims<ColMajor>()); + CALL_SUBTEST(test_multidims<RowMajor>()); + CALL_SUBTEST(test_holes<ColMajor>()); + CALL_SUBTEST(test_holes<RowMajor>()); + CALL_SUBTEST(test_full_redux<ColMajor>()); + CALL_SUBTEST(test_full_redux<RowMajor>()); + CALL_SUBTEST(test_contraction_of_contraction<ColMajor>()); + CALL_SUBTEST(test_contraction_of_contraction<RowMajor>()); + CALL_SUBTEST(test_expr<ColMajor>()); + CALL_SUBTEST(test_expr<RowMajor>()); + CALL_SUBTEST(test_out_of_order_contraction<ColMajor>()); + CALL_SUBTEST(test_out_of_order_contraction<RowMajor>()); + CALL_SUBTEST(test_consistency<ColMajor>()); + CALL_SUBTEST(test_consistency<RowMajor>()); + CALL_SUBTEST(test_large_contraction<ColMajor>()); + CALL_SUBTEST(test_large_contraction<RowMajor>()); + CALL_SUBTEST(test_matrix_vector<ColMajor>()); + CALL_SUBTEST(test_matrix_vector<RowMajor>()); + CALL_SUBTEST(test_tensor_vector<ColMajor>()); + CALL_SUBTEST(test_tensor_vector<RowMajor>()); + CALL_SUBTEST(test_small_blocking_factors<ColMajor>()); + CALL_SUBTEST(test_small_blocking_factors<RowMajor>()); } diff --git a/unsupported/test/cxx11_tensor_cuda.cpp b/unsupported/test/cxx11_tensor_cuda.cpp new file mode 100644 index 000000000..059d23de1 --- /dev/null +++ b/unsupported/test/cxx11_tensor_cuda.cpp @@ -0,0 +1,474 @@ +// This file is part of Eigen, a lightweight C++ template library +// for linear algebra. +// +// Copyright (C) 2014 Benoit Steiner <benoit.steiner.goog@gmail.com> +// +// 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/. + +// TODO(mdevin): Free the cuda memory. + +#define EIGEN_TEST_NO_LONGDOUBLE +#define EIGEN_TEST_NO_COMPLEX +#define EIGEN_TEST_FUNC cxx11_tensor_cuda +#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int +#define EIGEN_USE_GPU + + +#include "main.h" +#include <unsupported/Eigen/CXX11/Tensor> + +using Eigen::Tensor; + +void test_cuda_elementwise_small() { + Tensor<float, 1> in1(Eigen::array<int, 1>(2)); + Tensor<float, 1> in2(Eigen::array<int, 1>(2)); + Tensor<float, 1> out(Eigen::array<int, 1>(2)); + in1.setRandom(); + in2.setRandom(); + + std::size_t in1_bytes = in1.size() * sizeof(float); + std::size_t in2_bytes = in2.size() * sizeof(float); + std::size_t out_bytes = out.size() * sizeof(float); + + float* d_in1; + float* d_in2; + float* d_out; + cudaMalloc((void**)(&d_in1), in1_bytes); + cudaMalloc((void**)(&d_in2), in2_bytes); + cudaMalloc((void**)(&d_out), out_bytes); + + cudaMemcpy(d_in1, in1.data(), in1_bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_in2, in2.data(), in2_bytes, cudaMemcpyHostToDevice); + + cudaStream_t stream; + assert(cudaStreamCreate(&stream) == cudaSuccess); + Eigen::GpuDevice gpu_device(&stream); + + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_in1( + d_in1, Eigen::array<int, 1>(2)); + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_in2( + d_in2, Eigen::array<int, 1>(2)); + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_out( + d_out, Eigen::array<int, 1>(2)); + + gpu_out.device(gpu_device) = gpu_in1 + gpu_in2; + + assert(cudaMemcpyAsync(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost, + gpu_device.stream()) == cudaSuccess); + assert(cudaStreamSynchronize(gpu_device.stream()) == cudaSuccess); + + for (int i = 0; i < 2; ++i) { + VERIFY_IS_APPROX( + out(Eigen::array<int, 1>(i)), + in1(Eigen::array<int, 1>(i)) + in2(Eigen::array<int, 1>(i))); + } +} + +void test_cuda_elementwise() +{ + Tensor<float, 3> in1(Eigen::array<int, 3>(72,53,97)); + Tensor<float, 3> in2(Eigen::array<int, 3>(72,53,97)); + Tensor<float, 3> in3(Eigen::array<int, 3>(72,53,97)); + Tensor<float, 3> out(Eigen::array<int, 3>(72,53,97)); + in1.setRandom(); + in2.setRandom(); + in3.setRandom(); + + std::size_t in1_bytes = in1.size() * sizeof(float); + std::size_t in2_bytes = in2.size() * sizeof(float); + std::size_t in3_bytes = in3.size() * sizeof(float); + std::size_t out_bytes = out.size() * sizeof(float); + + float* d_in1; + float* d_in2; + float* d_in3; + float* d_out; + cudaMalloc((void**)(&d_in1), in1_bytes); + cudaMalloc((void**)(&d_in2), in2_bytes); + cudaMalloc((void**)(&d_in3), in3_bytes); + cudaMalloc((void**)(&d_out), out_bytes); + + cudaMemcpy(d_in1, in1.data(), in1_bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_in2, in2.data(), in2_bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_in3, in3.data(), in3_bytes, cudaMemcpyHostToDevice); + + cudaStream_t stream; + assert(cudaStreamCreate(&stream) == cudaSuccess); + Eigen::GpuDevice gpu_device(&stream); + + Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_in1(d_in1, Eigen::array<int, 3>(72,53,97)); + Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_in2(d_in2, Eigen::array<int, 3>(72,53,97)); + Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_in3(d_in3, Eigen::array<int, 3>(72,53,97)); + Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_out(d_out, Eigen::array<int, 3>(72,53,97)); + + gpu_out.device(gpu_device) = gpu_in1 + gpu_in2 * gpu_in3; + + assert(cudaMemcpyAsync(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost, gpu_device.stream()) == cudaSuccess); + assert(cudaStreamSynchronize(gpu_device.stream()) == cudaSuccess); + + for (int i = 0; i < 72; ++i) { + for (int j = 0; j < 53; ++j) { + for (int k = 0; k < 97; ++k) { + VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k)) * in3(Eigen::array<int, 3>(i,j,k))); + } + } + } +} + + +void test_cuda_reduction() +{ + Tensor<float, 4> in1(Eigen::array<int, 4>(72,53,97,113)); + Tensor<float, 2> out(Eigen::array<int, 2>(72,97)); + in1.setRandom(); + + std::size_t in1_bytes = in1.size() * sizeof(float); + std::size_t out_bytes = out.size() * sizeof(float); + + float* d_in1; + float* d_out; + cudaMalloc((void**)(&d_in1), in1_bytes); + cudaMalloc((void**)(&d_out), out_bytes); + + cudaMemcpy(d_in1, in1.data(), in1_bytes, cudaMemcpyHostToDevice); + + cudaStream_t stream; + assert(cudaStreamCreate(&stream) == cudaSuccess); + Eigen::GpuDevice gpu_device(&stream); + + Eigen::TensorMap<Eigen::Tensor<float, 4> > gpu_in1(d_in1, Eigen::array<int, 4>(72,53,97,113)); + Eigen::TensorMap<Eigen::Tensor<float, 2> > gpu_out(d_out, Eigen::array<int, 2>(72,97)); + + array<int, 2> reduction_axis; + reduction_axis[0] = 1; + reduction_axis[1] = 3; + + gpu_out.device(gpu_device) = gpu_in1.maximum(reduction_axis); + + assert(cudaMemcpyAsync(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost, gpu_device.stream()) == cudaSuccess); + assert(cudaStreamSynchronize(gpu_device.stream()) == cudaSuccess); + + for (int i = 0; i < 72; ++i) { + for (int j = 0; j < 97; ++j) { + float expected = 0; + for (int k = 0; k < 53; ++k) { + for (int l = 0; l < 113; ++l) { + expected = + std::max<float>(expected, in1(Eigen::array<int, 4>(i, k, j, l))); + } + } + VERIFY_IS_APPROX(out(Eigen::array<int, 2>(i,j)), expected); + } + } +} + +template<int DataLayout> +static void test_cuda_contraction() +{ + // with these dimensions, the output has 300 * 140 elements, which is + // more than 30 * 1024, which is the number of threads in blocks on + // a 15 SM GK110 GPU + Tensor<float, 4, DataLayout> t_left(Eigen::array<int, 4>(6, 50, 3, 31)); + Tensor<float, 5, DataLayout> t_right(Eigen::array<int, 5>(3, 31, 7, 20, 1)); + Tensor<float, 5, DataLayout> t_result(Eigen::array<int, 5>(6, 50, 7, 20, 1)); + + t_left.setRandom(); + t_right.setRandom(); + + std::size_t t_left_bytes = t_left.size() * sizeof(float); + std::size_t t_right_bytes = t_right.size() * sizeof(float); + std::size_t t_result_bytes = t_result.size() * sizeof(float); + + float* d_t_left; + float* d_t_right; + float* d_t_result; + + cudaMalloc((void**)(&d_t_left), t_left_bytes); + cudaMalloc((void**)(&d_t_right), t_right_bytes); + cudaMalloc((void**)(&d_t_result), t_result_bytes); + + cudaMemcpy(d_t_left, t_left.data(), t_left_bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_t_right, t_right.data(), t_right_bytes, cudaMemcpyHostToDevice); + + cudaStream_t stream; + assert(cudaStreamCreate(&stream) == cudaSuccess); + Eigen::GpuDevice gpu_device(&stream); + + Eigen::TensorMap<Eigen::Tensor<float, 4, DataLayout> > + gpu_t_left(d_t_left, Eigen::array<int, 4>(6, 50, 3, 31)); + Eigen::TensorMap<Eigen::Tensor<float, 5, DataLayout> > + gpu_t_right(d_t_right, Eigen::array<int, 5>(3, 31, 7, 20, 1)); + Eigen::TensorMap<Eigen::Tensor<float, 5, DataLayout> > + gpu_t_result(d_t_result, Eigen::array<int, 5>(6, 50, 7, 20, 1)); + + typedef Eigen::Map<Eigen::Matrix<float, Dynamic, Dynamic, DataLayout> > MapXf; + MapXf m_left(t_left.data(), 300, 93); + MapXf m_right(t_right.data(), 93, 140); + Eigen::Matrix<float, Dynamic, Dynamic, DataLayout> m_result(300, 140); + + typedef Tensor<float, 1>::DimensionPair DimPair; + Eigen::array<DimPair, 2> dims; + dims[0] = DimPair(2, 0); + dims[1] = DimPair(3, 1); + + m_result = m_left * m_right; + gpu_t_result.device(gpu_device) = gpu_t_left.contract(gpu_t_right, dims); + + cudaMemcpy(t_result.data(), d_t_result, t_result_bytes, cudaMemcpyDeviceToHost); + + for (size_t i = 0; i < t_result.dimensions().TotalSize(); i++) { + if (fabs(t_result.data()[i] - m_result.data()[i]) >= 1e-4) { + cout << "mismatch detected at index " << i << ": " << t_result.data()[i] << " vs " << m_result.data()[i] << endl; + assert(false); + } + } +} + +static void test_cuda_convolution_1d() +{ + Tensor<float, 4> input(Eigen::array<int, 4>(74,37,11,137)); + Tensor<float, 1> kernel(Eigen::array<int, 1>(4)); + Tensor<float, 4> out(Eigen::array<int, 4>(74,34,11,137)); + input = input.constant(10.0f) + input.random(); + kernel = kernel.constant(7.0f) + kernel.random(); + + std::size_t input_bytes = input.size() * sizeof(float); + std::size_t kernel_bytes = kernel.size() * sizeof(float); + std::size_t out_bytes = out.size() * sizeof(float); + + float* d_input; + float* d_kernel; + float* d_out; + cudaMalloc((void**)(&d_input), input_bytes); + cudaMalloc((void**)(&d_kernel), kernel_bytes); + cudaMalloc((void**)(&d_out), out_bytes); + + cudaMemcpy(d_input, input.data(), input_bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_kernel, kernel.data(), kernel_bytes, cudaMemcpyHostToDevice); + + cudaStream_t stream; + assert(cudaStreamCreate(&stream) == cudaSuccess); + Eigen::GpuDevice gpu_device(&stream); + + Eigen::TensorMap<Eigen::Tensor<float, 4> > gpu_input(d_input, Eigen::array<int, 4>(74,37,11,137)); + Eigen::TensorMap<Eigen::Tensor<float, 1> > gpu_kernel(d_kernel, Eigen::array<int, 1>(4)); + Eigen::TensorMap<Eigen::Tensor<float, 4> > gpu_out(d_out, Eigen::array<int, 4>(74,34,11,137)); + + Eigen::array<int, 1> dims(1); + gpu_out.device(gpu_device) = gpu_input.convolve(gpu_kernel, dims); + + assert(cudaMemcpyAsync(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost, gpu_device.stream()) == cudaSuccess); + assert(cudaStreamSynchronize(gpu_device.stream()) == cudaSuccess); + + for (int i = 0; i < 74; ++i) { + for (int j = 0; j < 34; ++j) { + for (int k = 0; k < 11; ++k) { + for (int l = 0; l < 137; ++l) { + const float result = out(Eigen::array<int, 4>(i,j,k,l)); + const float expected = input(Eigen::array<int, 4>(i,j+0,k,l)) * kernel(Eigen::array<int, 1>(0)) + + input(Eigen::array<int, 4>(i,j+1,k,l)) * kernel(Eigen::array<int, 1>(1)) + + input(Eigen::array<int, 4>(i,j+2,k,l)) * kernel(Eigen::array<int, 1>(2)) + + input(Eigen::array<int, 4>(i,j+3,k,l)) * kernel(Eigen::array<int, 1>(3)); + VERIFY_IS_APPROX(result, expected); + } + } + } + } +} + + +static void test_cuda_convolution_2d() +{ + Tensor<float, 4> input(Eigen::array<int, 4>(74,37,11,137)); + Tensor<float, 2> kernel(Eigen::array<int, 2>(3,4)); + Tensor<float, 4> out(Eigen::array<int, 4>(74,35,8,137)); + input = input.constant(10.0f) + input.random(); + kernel = kernel.constant(7.0f) + kernel.random(); + + std::size_t input_bytes = input.size() * sizeof(float); + std::size_t kernel_bytes = kernel.size() * sizeof(float); + std::size_t out_bytes = out.size() * sizeof(float); + + float* d_input; + float* d_kernel; + float* d_out; + cudaMalloc((void**)(&d_input), input_bytes); + cudaMalloc((void**)(&d_kernel), kernel_bytes); + cudaMalloc((void**)(&d_out), out_bytes); + + cudaMemcpy(d_input, input.data(), input_bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_kernel, kernel.data(), kernel_bytes, cudaMemcpyHostToDevice); + + cudaStream_t stream; + assert(cudaStreamCreate(&stream) == cudaSuccess); + Eigen::GpuDevice gpu_device(&stream); + + Eigen::TensorMap<Eigen::Tensor<float, 4> > gpu_input(d_input, Eigen::array<int, 4>(74,37,11,137)); + Eigen::TensorMap<Eigen::Tensor<float, 2> > gpu_kernel(d_kernel, Eigen::array<int, 2>(3,4)); + Eigen::TensorMap<Eigen::Tensor<float, 4> > gpu_out(d_out, Eigen::array<int, 4>(74,35,8,137)); + + Eigen::array<int, 2> dims(1,2); + gpu_out.device(gpu_device) = gpu_input.convolve(gpu_kernel, dims); + + assert(cudaMemcpyAsync(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost, gpu_device.stream()) == cudaSuccess); + assert(cudaStreamSynchronize(gpu_device.stream()) == cudaSuccess); + + for (int i = 0; i < 74; ++i) { + for (int j = 0; j < 35; ++j) { + for (int k = 0; k < 8; ++k) { + for (int l = 0; l < 137; ++l) { + const float result = out(Eigen::array<int, 4>(i,j,k,l)); + const float expected = input(Eigen::array<int, 4>(i,j+0,k+0,l)) * kernel(Eigen::array<int, 2>(0,0)) + + input(Eigen::array<int, 4>(i,j+1,k+0,l)) * kernel(Eigen::array<int, 2>(1,0)) + + input(Eigen::array<int, 4>(i,j+2,k+0,l)) * kernel(Eigen::array<int, 2>(2,0)) + + input(Eigen::array<int, 4>(i,j+0,k+1,l)) * kernel(Eigen::array<int, 2>(0,1)) + + input(Eigen::array<int, 4>(i,j+1,k+1,l)) * kernel(Eigen::array<int, 2>(1,1)) + + input(Eigen::array<int, 4>(i,j+2,k+1,l)) * kernel(Eigen::array<int, 2>(2,1)) + + input(Eigen::array<int, 4>(i,j+0,k+2,l)) * kernel(Eigen::array<int, 2>(0,2)) + + input(Eigen::array<int, 4>(i,j+1,k+2,l)) * kernel(Eigen::array<int, 2>(1,2)) + + input(Eigen::array<int, 4>(i,j+2,k+2,l)) * kernel(Eigen::array<int, 2>(2,2)) + + input(Eigen::array<int, 4>(i,j+0,k+3,l)) * kernel(Eigen::array<int, 2>(0,3)) + + input(Eigen::array<int, 4>(i,j+1,k+3,l)) * kernel(Eigen::array<int, 2>(1,3)) + + input(Eigen::array<int, 4>(i,j+2,k+3,l)) * kernel(Eigen::array<int, 2>(2,3)); + VERIFY_IS_APPROX(result, expected); + } + } + } + } +} + + +static void test_cuda_convolution_3d() +{ + Tensor<float, 5> input(Eigen::array<int, 5>(74,37,11,137,17)); + Tensor<float, 3> kernel(Eigen::array<int, 3>(3,4,2)); + Tensor<float, 5> out(Eigen::array<int, 5>(74,35,8,136,17)); + input = input.constant(10.0f) + input.random(); + kernel = kernel.constant(7.0f) + kernel.random(); + + std::size_t input_bytes = input.size() * sizeof(float); + std::size_t kernel_bytes = kernel.size() * sizeof(float); + std::size_t out_bytes = out.size() * sizeof(float); + + float* d_input; + float* d_kernel; + float* d_out; + cudaMalloc((void**)(&d_input), input_bytes); + cudaMalloc((void**)(&d_kernel), kernel_bytes); + cudaMalloc((void**)(&d_out), out_bytes); + + cudaMemcpy(d_input, input.data(), input_bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_kernel, kernel.data(), kernel_bytes, cudaMemcpyHostToDevice); + + cudaStream_t stream; + assert(cudaStreamCreate(&stream) == cudaSuccess); + Eigen::GpuDevice gpu_device(&stream); + + Eigen::TensorMap<Eigen::Tensor<float, 5> > gpu_input(d_input, Eigen::array<int, 5>(74,37,11,137,17)); + Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_kernel(d_kernel, Eigen::array<int, 3>(3,4,2)); + Eigen::TensorMap<Eigen::Tensor<float, 5> > gpu_out(d_out, Eigen::array<int, 5>(74,35,8,136,17)); + + Eigen::array<int, 3> dims(1,2,3); + gpu_out.device(gpu_device) = gpu_input.convolve(gpu_kernel, dims); + + assert(cudaMemcpyAsync(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost, gpu_device.stream()) == cudaSuccess); + assert(cudaStreamSynchronize(gpu_device.stream()) == cudaSuccess); + + for (int i = 0; i < 74; ++i) { + for (int j = 0; j < 35; ++j) { + for (int k = 0; k < 8; ++k) { + for (int l = 0; l < 136; ++l) { + for (int m = 0; m < 17; ++m) { + const float result = out(Eigen::array<int, 5>(i,j,k,l,m)); + const float expected = input(Eigen::array<int, 5>(i,j+0,k+0,l+0,m)) * kernel(Eigen::array<int, 3>(0,0,0)) + + input(Eigen::array<int, 5>(i,j+1,k+0,l+0,m)) * kernel(Eigen::array<int, 3>(1,0,0)) + + input(Eigen::array<int, 5>(i,j+2,k+0,l+0,m)) * kernel(Eigen::array<int, 3>(2,0,0)) + + input(Eigen::array<int, 5>(i,j+0,k+1,l+0,m)) * kernel(Eigen::array<int, 3>(0,1,0)) + + input(Eigen::array<int, 5>(i,j+1,k+1,l+0,m)) * kernel(Eigen::array<int, 3>(1,1,0)) + + input(Eigen::array<int, 5>(i,j+2,k+1,l+0,m)) * kernel(Eigen::array<int, 3>(2,1,0)) + + input(Eigen::array<int, 5>(i,j+0,k+2,l+0,m)) * kernel(Eigen::array<int, 3>(0,2,0)) + + input(Eigen::array<int, 5>(i,j+1,k+2,l+0,m)) * kernel(Eigen::array<int, 3>(1,2,0)) + + input(Eigen::array<int, 5>(i,j+2,k+2,l+0,m)) * kernel(Eigen::array<int, 3>(2,2,0)) + + input(Eigen::array<int, 5>(i,j+0,k+3,l+0,m)) * kernel(Eigen::array<int, 3>(0,3,0)) + + input(Eigen::array<int, 5>(i,j+1,k+3,l+0,m)) * kernel(Eigen::array<int, 3>(1,3,0)) + + input(Eigen::array<int, 5>(i,j+2,k+3,l+0,m)) * kernel(Eigen::array<int, 3>(2,3,0)) + + input(Eigen::array<int, 5>(i,j+0,k+0,l+1,m)) * kernel(Eigen::array<int, 3>(0,0,1)) + + input(Eigen::array<int, 5>(i,j+1,k+0,l+1,m)) * kernel(Eigen::array<int, 3>(1,0,1)) + + input(Eigen::array<int, 5>(i,j+2,k+0,l+1,m)) * kernel(Eigen::array<int, 3>(2,0,1)) + + input(Eigen::array<int, 5>(i,j+0,k+1,l+1,m)) * kernel(Eigen::array<int, 3>(0,1,1)) + + input(Eigen::array<int, 5>(i,j+1,k+1,l+1,m)) * kernel(Eigen::array<int, 3>(1,1,1)) + + input(Eigen::array<int, 5>(i,j+2,k+1,l+1,m)) * kernel(Eigen::array<int, 3>(2,1,1)) + + input(Eigen::array<int, 5>(i,j+0,k+2,l+1,m)) * kernel(Eigen::array<int, 3>(0,2,1)) + + input(Eigen::array<int, 5>(i,j+1,k+2,l+1,m)) * kernel(Eigen::array<int, 3>(1,2,1)) + + input(Eigen::array<int, 5>(i,j+2,k+2,l+1,m)) * kernel(Eigen::array<int, 3>(2,2,1)) + + input(Eigen::array<int, 5>(i,j+0,k+3,l+1,m)) * kernel(Eigen::array<int, 3>(0,3,1)) + + input(Eigen::array<int, 5>(i,j+1,k+3,l+1,m)) * kernel(Eigen::array<int, 3>(1,3,1)) + + input(Eigen::array<int, 5>(i,j+2,k+3,l+1,m)) * kernel(Eigen::array<int, 3>(2,3,1)); + VERIFY_IS_APPROX(result, expected); + } + } + } + } + } +} + +static float* CudaCopyFloat(float* data, int size) { + const int nbytes = size * sizeof(float); + float* result = NULL; + if (cudaMalloc((void**)(&result), nbytes) != cudaSuccess) { + return NULL; + } else { + if (data != NULL) { + cudaMemcpy(result, data, nbytes, cudaMemcpyHostToDevice); + } + return result; + } +} + +static void test_cuda_constant_broadcast() +{ + cudaStream_t stream; + assert(cudaStreamCreate(&stream) == cudaSuccess); + Eigen::GpuDevice gpu_device(&stream); + + Tensor<float, 1> t1(10); + for (int i = 0; i < 10; ++i) { + t1(i) = 10.0f * i; + } + float* t1_cuda = CudaCopyFloat(t1.data(), t1.size()); + Eigen::TensorMap<Eigen::Tensor<float, 1> > t1_gpu(t1_cuda, 10); + + Tensor<float, 1> t2(1); + t2 = t2.constant(20.0f); + float* t2_cuda = CudaCopyFloat(t2.data(), t2.size()); + Eigen::TensorMap<Eigen::TensorFixedSize<float, Sizes<1> > > t2_gpu(t2_cuda, 1); + + float* t3_cuda = CudaCopyFloat(NULL, 10); + Eigen::TensorMap<Eigen::Tensor<float, 1> > t3_gpu(t3_cuda, 10); + + t3_gpu.device(gpu_device) = + t1_gpu + t2_gpu.broadcast(Eigen::array<int, 1>(10)); + + Eigen::Tensor<float, 1> t3(10); + cudaMemcpy(t3.data(), t3_gpu.data(), 10 * sizeof(float), + cudaMemcpyDeviceToHost); + + for (int i = 0; i < 10; ++i) { + VERIFY_IS_APPROX(t3(i), t1(i) + t2(0)); + } +} + +void test_cxx11_tensor_cuda() +{ + CALL_SUBTEST(test_cuda_elementwise_small()); + CALL_SUBTEST(test_cuda_elementwise()); + CALL_SUBTEST(test_cuda_reduction()); + CALL_SUBTEST(test_cuda_contraction<ColMajor>()); + CALL_SUBTEST(test_cuda_contraction<RowMajor>()); + CALL_SUBTEST(test_cuda_convolution_1d()); + CALL_SUBTEST(test_cuda_convolution_2d()); + CALL_SUBTEST(test_cuda_convolution_3d()); + CALL_SUBTEST(test_cuda_constant_broadcast()); +} diff --git a/unsupported/test/cxx11_tensor_device.cpp b/unsupported/test/cxx11_tensor_device.cpp index 26465ee11..f2d7e4ce6 100644 --- a/unsupported/test/cxx11_tensor_device.cpp +++ b/unsupported/test/cxx11_tensor_device.cpp @@ -22,23 +22,23 @@ using Eigen::RowMajor; // Context for evaluation on cpu struct CPUContext { - CPUContext(const Eigen::Tensor<float, 3>& in1, Eigen::Tensor<float, 3>& in2, Eigen::Tensor<float, 3>& out) : in1_(in1), in2_(in2), out_(out), kernel_1d_(2), kernel_2d_(Eigen::array<int, 2>(2,2)), kernel_3d_(Eigen::array<int, 3>(2,2,2)) { + CPUContext(const Eigen::Tensor<float, 3>& in1, Eigen::Tensor<float, 3>& in2, Eigen::Tensor<float, 3>& out) : in1_(in1), in2_(in2), out_(out), kernel_1d_(2), kernel_2d_(2,2), kernel_3d_(2,2,2) { kernel_1d_(0) = 3.14f; kernel_1d_(1) = 2.7f; - kernel_2d_(Eigen::array<int, 2>(0,0)) = 3.14f; - kernel_2d_(Eigen::array<int, 2>(1,0)) = 2.7f; - kernel_2d_(Eigen::array<int, 2>(0,1)) = 0.2f; - kernel_2d_(Eigen::array<int, 2>(1,1)) = 7.0f; - - kernel_3d_(Eigen::array<int, 3>(0,0,0)) = 3.14f; - kernel_3d_(Eigen::array<int, 3>(0,1,0)) = 2.7f; - kernel_3d_(Eigen::array<int, 3>(0,0,1)) = 0.2f; - kernel_3d_(Eigen::array<int, 3>(0,1,1)) = 7.0f; - kernel_3d_(Eigen::array<int, 3>(1,0,0)) = -1.0f; - kernel_3d_(Eigen::array<int, 3>(1,1,0)) = -0.3f; - kernel_3d_(Eigen::array<int, 3>(1,0,1)) = -0.7f; - kernel_3d_(Eigen::array<int, 3>(1,1,1)) = -0.5f; + kernel_2d_(0,0) = 3.14f; + kernel_2d_(1,0) = 2.7f; + kernel_2d_(0,1) = 0.2f; + kernel_2d_(1,1) = 7.0f; + + kernel_3d_(0,0,0) = 3.14f; + kernel_3d_(0,1,0) = 2.7f; + kernel_3d_(0,0,1) = 0.2f; + kernel_3d_(0,1,1) = 7.0f; + kernel_3d_(1,0,0) = -1.0f; + kernel_3d_(1,1,0) = -0.3f; + kernel_3d_(1,0,1) = -0.7f; + kernel_3d_(1,1,1) = -0.5f; } const Eigen::DefaultDevice& device() const { return cpu_device_; } @@ -93,8 +93,8 @@ struct GPUContext { const Eigen::TensorMap<Eigen::Tensor<float, 3> >& in2() const { return in2_; } Eigen::TensorMap<Eigen::Tensor<float, 3> >& out() { return out_; } Eigen::TensorMap<Eigen::Tensor<float, 1> > kernel1d() const { return Eigen::TensorMap<Eigen::Tensor<float, 1> >(kernel_1d_, 2); } - Eigen::TensorMap<Eigen::Tensor<float, 2> > kernel2d() const { return Eigen::TensorMap<Eigen::Tensor<float, 2> >(kernel_2d_, Eigen::array<int, 2>(2, 2)); } - Eigen::TensorMap<Eigen::Tensor<float, 3> > kernel3d() const { return Eigen::TensorMap<Eigen::Tensor<float, 3> >(kernel_3d_, Eigen::array<int, 3>(2, 2, 2)); } + Eigen::TensorMap<Eigen::Tensor<float, 2> > kernel2d() const { return Eigen::TensorMap<Eigen::Tensor<float, 2> >(kernel_2d_, 2, 2); } + Eigen::TensorMap<Eigen::Tensor<float, 3> > kernel3d() const { return Eigen::TensorMap<Eigen::Tensor<float, 3> >(kernel_3d_, 2, 2, 2); } private: const Eigen::TensorMap<Eigen::Tensor<float, 3> >& in1_; @@ -150,8 +150,8 @@ static void test_contraction(Context* context) template <typename Context> static void test_1d_convolution(Context* context) { - Eigen::DSizes<int, 3> indices(Eigen::array<int, 3>(0,0,0)); - Eigen::DSizes<int, 3> sizes(Eigen::array<int, 3>(40,49,70)); + Eigen::DSizes<int, 3> indices(0,0,0); + Eigen::DSizes<int, 3> sizes(40,49,70); Eigen::array<int, 1> dims(1); context->out().slice(indices, sizes).device(context->device()) = context->in1().convolve(context->kernel1d(), dims); @@ -160,8 +160,8 @@ static void test_1d_convolution(Context* context) template <typename Context> static void test_2d_convolution(Context* context) { - Eigen::DSizes<int, 3> indices(Eigen::array<int, 3>(0,0,0)); - Eigen::DSizes<int, 3> sizes(Eigen::array<int, 3>(40,49,69)); + Eigen::DSizes<int, 3> indices(0,0,0); + Eigen::DSizes<int, 3> sizes(40,49,69); Eigen::array<int, 2> dims(1,2); context->out().slice(indices, sizes).device(context->device()) = context->in1().convolve(context->kernel2d(), dims); @@ -170,8 +170,8 @@ static void test_2d_convolution(Context* context) template <typename Context> static void test_3d_convolution(Context* context) { - Eigen::DSizes<int, 3> indices(Eigen::array<int, 3>(0,0,0)); - Eigen::DSizes<int, 3> sizes(Eigen::array<int, 3>(39,49,69)); + Eigen::DSizes<int, 3> indices(0,0,0); + Eigen::DSizes<int, 3> sizes(39,49,69); Eigen::array<int, 3> dims(0,1,2); context->out().slice(indices, sizes).device(context->device()) = context->in1().convolve(context->kernel3d(), dims); @@ -179,9 +179,9 @@ static void test_3d_convolution(Context* context) static void test_cpu() { - Eigen::Tensor<float, 3> in1(Eigen::array<int, 3>(40,50,70)); - Eigen::Tensor<float, 3> in2(Eigen::array<int, 3>(40,50,70)); - Eigen::Tensor<float, 3> out(Eigen::array<int, 3>(40,50,70)); + Eigen::Tensor<float, 3> in1(40,50,70); + Eigen::Tensor<float, 3> in2(40,50,70); + Eigen::Tensor<float, 3> out(40,50,70); in1 = in1.random() + in1.constant(10.0f); in2 = in2.random() + in2.constant(10.0f); @@ -191,7 +191,7 @@ static void test_cpu() { for (int i = 0; i < 40; ++i) { for (int j = 0; j < 50; ++j) { for (int k = 0; k < 70; ++k) { - VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k)) * 3.14f + 2.718f); + VERIFY_IS_APPROX(out(i,j,k), in1(i,j,k) + in2(i,j,k) * 3.14f + 2.718f); } } } @@ -200,7 +200,7 @@ static void test_cpu() { for (int i = 0; i < 40; ++i) { for (int j = 0; j < 50; ++j) { for (int k = 0; k < 70; ++k) { - VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), (in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k))) * 3.14f + 2.718f); + VERIFY_IS_APPROX(out(i,j,k), (in1(i,j,k) + in2(i,j,k)) * 3.14f + 2.718f); } } } @@ -209,7 +209,7 @@ static void test_cpu() { for (int i = 0; i < 40; ++i) { for (int j = 0; j < 50; ++j) { for (int k = 0; k < 70; ++k) { - VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k)) * 3.14f + 2.718f); + VERIFY_IS_APPROX(out(i,j,k), in1(i,j,k) + in2(i,j,k) * 3.14f + 2.718f); } } } @@ -217,11 +217,11 @@ static void test_cpu() { test_contraction(&context); for (int i = 0; i < 40; ++i) { for (int j = 0; j < 40; ++j) { - const float result = out(Eigen::array<int, 3>(i,j,0)); + const float result = out(i,j,0); float expected = 0; for (int k = 0; k < 50; ++k) { for (int l = 0; l < 70; ++l) { - expected += in1(Eigen::array<int, 3>(i, k, l)) * in2(Eigen::array<int, 3>(j, k, l)); + expected += in1(i, k, l) * in2(j, k, l); } } VERIFY_IS_APPROX(expected, result); @@ -232,7 +232,7 @@ static void test_cpu() { for (int i = 0; i < 40; ++i) { for (int j = 0; j < 49; ++j) { for (int k = 0; k < 70; ++k) { - VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f)); + VERIFY_IS_APPROX(out(i,j,k), (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f)); } } } @@ -241,9 +241,9 @@ static void test_cpu() { for (int i = 0; i < 40; ++i) { for (int j = 0; j < 49; ++j) { for (int k = 0; k < 69; ++k) { - const float result = out(Eigen::array<int, 3>(i,j,k)); - const float expected = (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f) + - (in1(Eigen::array<int, 3>(i,j,k+1)) * 0.2f + in1(Eigen::array<int, 3>(i,j+1,k+1)) * 7.0f); + const float result = out(i,j,k); + const float expected = (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f) + + (in1(i,j,k+1) * 0.2f + in1(i,j+1,k+1) * 7.0f); if (fabs(expected) < 1e-4 && fabs(result) < 1e-4) { continue; } @@ -256,11 +256,11 @@ static void test_cpu() { for (int i = 0; i < 39; ++i) { for (int j = 0; j < 49; ++j) { for (int k = 0; k < 69; ++k) { - const float result = out(Eigen::array<int, 3>(i,j,k)); - const float expected = (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f + - in1(Eigen::array<int, 3>(i,j,k+1)) * 0.2f + in1(Eigen::array<int, 3>(i,j+1,k+1)) * 7.0f) + - (in1(Eigen::array<int, 3>(i+1,j,k)) * -1.0f + in1(Eigen::array<int, 3>(i+1,j+1,k)) * -0.3f + - in1(Eigen::array<int, 3>(i+1,j,k+1)) * -0.7f + in1(Eigen::array<int, 3>(i+1,j+1,k+1)) * -0.5f); + const float result = out(i,j,k); + const float expected = (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f + + in1(i,j,k+1) * 0.2f + in1(i,j+1,k+1) * 7.0f) + + (in1(i+1,j,k) * -1.0f + in1(i+1,j+1,k) * -0.3f + + in1(i+1,j,k+1) * -0.7f + in1(i+1,j+1,k+1) * -0.5f); if (fabs(expected) < 1e-4 && fabs(result) < 1e-4) { continue; } @@ -271,9 +271,9 @@ static void test_cpu() { } static void test_gpu() { - Eigen::Tensor<float, 3> in1(Eigen::array<int, 3>(40,50,70)); - Eigen::Tensor<float, 3> in2(Eigen::array<int, 3>(40,50,70)); - Eigen::Tensor<float, 3> out(Eigen::array<int, 3>(40,50,70)); + Eigen::Tensor<float, 3> in1(40,50,70); + Eigen::Tensor<float, 3> in2(40,50,70); + Eigen::Tensor<float, 3> out(40,50,70); in1 = in1.random() + in1.constant(10.0f); in2 = in2.random() + in2.constant(10.0f); @@ -291,9 +291,9 @@ static void test_gpu() { cudaMemcpy(d_in1, in1.data(), in1_bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_in2, in2.data(), in2_bytes, cudaMemcpyHostToDevice); - Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_in1(d_in1, Eigen::array<int, 3>(40,50,70)); - Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_in2(d_in2, Eigen::array<int, 3>(40,50,70)); - Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_out(d_out, Eigen::array<int, 3>(40,50,70)); + Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_in1(d_in1, 40,50,70); + Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_in2(d_in2, 40,50,70); + Eigen::TensorMap<Eigen::Tensor<float, 3> > gpu_out(d_out, 40,50,70); GPUContext context(gpu_in1, gpu_in2, gpu_out); test_contextual_eval(&context); @@ -301,7 +301,7 @@ static void test_gpu() { for (int i = 0; i < 40; ++i) { for (int j = 0; j < 50; ++j) { for (int k = 0; k < 70; ++k) { - VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k)) * 3.14f + 2.718f); + VERIFY_IS_APPROX(out(i,j,k), in1(i,j,k) + in2(i,j,k) * 3.14f + 2.718f); } } } @@ -311,7 +311,7 @@ static void test_gpu() { for (int i = 0; i < 40; ++i) { for (int j = 0; j < 50; ++j) { for (int k = 0; k < 70; ++k) { - VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), (in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k))) * 3.14f + 2.718f); + VERIFY_IS_APPROX(out(i,j,k), (in1(i,j,k) + in2(i,j,k)) * 3.14f + 2.718f); } } } @@ -321,7 +321,7 @@ static void test_gpu() { for (int i = 0; i < 40; ++i) { for (int j = 0; j < 50; ++j) { for (int k = 0; k < 70; ++k) { - VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), in1(Eigen::array<int, 3>(i,j,k)) + in2(Eigen::array<int, 3>(i,j,k)) * 3.14f + 2.718f); + VERIFY_IS_APPROX(out(i,j,k), in1(i,j,k) + in2(i,j,k) * 3.14f + 2.718f); } } } @@ -330,11 +330,11 @@ static void test_gpu() { assert(cudaMemcpy(out.data(), d_out, out_bytes, cudaMemcpyDeviceToHost) == cudaSuccess); for (int i = 0; i < 40; ++i) { for (int j = 0; j < 40; ++j) { - const float result = out(Eigen::array<int, 3>(i,j,0)); + const float result = out(i,j,0); float expected = 0; for (int k = 0; k < 50; ++k) { for (int l = 0; l < 70; ++l) { - expected += in1(Eigen::array<int, 3>(i, k, l)) * in2(Eigen::array<int, 3>(j, k, l)); + expected += in1(i, k, l) * in2(j, k, l); } } VERIFY_IS_APPROX(expected, result); @@ -347,7 +347,7 @@ static void test_gpu() { for (int i = 0; i < 40; ++i) { for (int j = 0; j < 49; ++j) { for (int k = 0; k < 70; ++k) { - VERIFY_IS_APPROX(out(Eigen::array<int, 3>(i,j,k)), (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f)); + VERIFY_IS_APPROX(out(i,j,k), (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f)); } } } @@ -358,9 +358,9 @@ static void test_gpu() { for (int i = 0; i < 40; ++i) { for (int j = 0; j < 49; ++j) { for (int k = 0; k < 69; ++k) { - const float result = out(Eigen::array<int, 3>(i,j,k)); - const float expected = (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f + - in1(Eigen::array<int, 3>(i,j,k+1)) * 0.2f + in1(Eigen::array<int, 3>(i,j+1,k+1)) * 7.0f); + const float result = out(i,j,k); + const float expected = (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f + + in1(i,j,k+1) * 0.2f + in1(i,j+1,k+1) * 7.0f); VERIFY_IS_APPROX(expected, result); } } @@ -372,11 +372,11 @@ static void test_gpu() { for (int i = 0; i < 39; ++i) { for (int j = 0; j < 49; ++j) { for (int k = 0; k < 69; ++k) { - const float result = out(Eigen::array<int, 3>(i,j,k)); - const float expected = (in1(Eigen::array<int, 3>(i,j,k)) * 3.14f + in1(Eigen::array<int, 3>(i,j+1,k)) * 2.7f + - in1(Eigen::array<int, 3>(i,j,k+1)) * 0.2f + in1(Eigen::array<int, 3>(i,j+1,k+1)) * 7.0f + - in1(Eigen::array<int, 3>(i+1,j,k)) * -1.0f + in1(Eigen::array<int, 3>(i+1,j+1,k)) * -0.3f + - in1(Eigen::array<int, 3>(i+1,j,k+1)) * -0.7f + in1(Eigen::array<int, 3>(i+1,j+1,k+1)) * -0.5f); + const float result = out(i,j,k); + const float expected = (in1(i,j,k) * 3.14f + in1(i,j+1,k) * 2.7f + + in1(i,j,k+1) * 0.2f + in1(i,j+1,k+1) * 7.0f + + in1(i+1,j,k) * -1.0f + in1(i+1,j+1,k) * -0.3f + + in1(i+1,j,k+1) * -0.7f + in1(i+1,j+1,k+1) * -0.5f); VERIFY_IS_APPROX(expected, result); } } diff --git a/unsupported/test/cxx11_tensor_dimension.cpp b/unsupported/test/cxx11_tensor_dimension.cpp index c806b623f..0cc4e86f7 100644 --- a/unsupported/test/cxx11_tensor_dimension.cpp +++ b/unsupported/test/cxx11_tensor_dimension.cpp @@ -16,12 +16,15 @@ using Eigen::Tensor; static void test_dynamic_size() { - Eigen::DSizes<int, 3> dimensions(Eigen::array<int, 3>{{2,3,7}}); + Eigen::DSizes<int, 3> dimensions(2,3,7); VERIFY_IS_EQUAL((int)Eigen::internal::array_get<0>(dimensions), 2); VERIFY_IS_EQUAL((int)Eigen::internal::array_get<1>(dimensions), 3); VERIFY_IS_EQUAL((int)Eigen::internal::array_get<2>(dimensions), 7); VERIFY_IS_EQUAL(dimensions.TotalSize(), (size_t)2*3*7); + VERIFY_IS_EQUAL((int)dimensions[0], 2); + VERIFY_IS_EQUAL((int)dimensions[1], 3); + VERIFY_IS_EQUAL((int)dimensions[2], 7); } static void test_fixed_size() @@ -37,9 +40,9 @@ static void test_fixed_size() static void test_match() { - Eigen::DSizes<int, 3> dyn(Eigen::array<int, 3>{{2,3,7}}); + Eigen::DSizes<int, 3> dyn(2,3,7); Eigen::Sizes<2,3,7> stat; - VERIFY_IS_EQUAL(Eigen::internal::dimensions_match(dyn, stat), true); + VERIFY_IS_EQUAL(Eigen::dimensions_match(dyn, stat), true); } diff --git a/unsupported/test/cxx11_tensor_expr.cpp b/unsupported/test/cxx11_tensor_expr.cpp index e85fcbfa9..792fdeade 100644 --- a/unsupported/test/cxx11_tensor_expr.cpp +++ b/unsupported/test/cxx11_tensor_expr.cpp @@ -125,6 +125,12 @@ static void test_3d() mat7 = mat1.cwiseMax(mat5 * 2.0f).exp(); Tensor<float, 3, RowMajor> mat8(2,3,7); mat8 = (-mat2).exp() * 3.14f; + Tensor<float, 3, RowMajor> mat9(2,3,7); + mat9 = mat2 + 3.14f; + Tensor<float, 3, RowMajor> mat10(2,3,7); + mat10 = mat2 - 3.14f; + Tensor<float, 3, RowMajor> mat11(2,3,7); + mat11 = mat2 / 3.14f; val = 1.0; for (int i = 0; i < 2; ++i) { @@ -136,6 +142,9 @@ static void test_3d() VERIFY_IS_APPROX(mat6(i,j,k), sqrtf(val) * 3.14f); VERIFY_IS_APPROX(mat7(i,j,k), expf((std::max)(val, mat5(i,j,k) * 2.0f))); VERIFY_IS_APPROX(mat8(i,j,k), expf(-val) * 3.14f); + VERIFY_IS_APPROX(mat9(i,j,k), val + 3.14f); + VERIFY_IS_APPROX(mat10(i,j,k), val - 3.14f); + VERIFY_IS_APPROX(mat11(i,j,k), val / 3.14f); val += 1.0; } } @@ -172,6 +181,36 @@ static void test_constants() } } +static void test_boolean() +{ + Tensor<int, 1> vec(6); + std::copy_n(std::begin({0, 1, 2, 3, 4, 5}), 6, vec.data()); + + // Test ||. + Tensor<bool, 1> bool1 = vec < vec.constant(1) || vec > vec.constant(4); + VERIFY_IS_EQUAL(bool1[0], true); + VERIFY_IS_EQUAL(bool1[1], false); + VERIFY_IS_EQUAL(bool1[2], false); + VERIFY_IS_EQUAL(bool1[3], false); + VERIFY_IS_EQUAL(bool1[4], false); + VERIFY_IS_EQUAL(bool1[5], true); + + // Test &&, including cast of operand vec. + Tensor<bool, 1> bool2 = vec.cast<bool>() && vec < vec.constant(4); + VERIFY_IS_EQUAL(bool2[0], false); + VERIFY_IS_EQUAL(bool2[1], true); + VERIFY_IS_EQUAL(bool2[2], true); + VERIFY_IS_EQUAL(bool2[3], true); + VERIFY_IS_EQUAL(bool2[4], false); + VERIFY_IS_EQUAL(bool2[5], false); + + // Compilation tests: + // Test Tensor<bool> against results of cast or comparison; verifies that + // CoeffReturnType is set to match Op return type of bool for Unary and Binary + // Ops. + Tensor<bool, 1> bool3 = vec.cast<bool>() && bool2; + bool3 = vec < vec.constant(4) && bool2; +} static void test_functors() { @@ -258,6 +297,7 @@ void test_cxx11_tensor_expr() CALL_SUBTEST(test_2d()); CALL_SUBTEST(test_3d()); CALL_SUBTEST(test_constants()); + CALL_SUBTEST(test_boolean()); CALL_SUBTEST(test_functors()); CALL_SUBTEST(test_type_casting()); CALL_SUBTEST(test_select()); diff --git a/unsupported/test/cxx11_tensor_forced_eval.cpp b/unsupported/test/cxx11_tensor_forced_eval.cpp index 529584a7b..ad9de867d 100644 --- a/unsupported/test/cxx11_tensor_forced_eval.cpp +++ b/unsupported/test/cxx11_tensor_forced_eval.cpp @@ -45,7 +45,34 @@ static void test_simple() } +static void test_const() +{ + MatrixXf input(3,3); + input.setRandom(); + MatrixXf output = input; + output.rowwise() -= input.colwise().maxCoeff(); + + Eigen::array<int, 1> depth_dim; + depth_dim[0] = 0; + Tensor<float, 2>::Dimensions dims2d; + dims2d[0] = 1; + dims2d[1] = 3; + Eigen::array<int, 2> bcast; + bcast[0] = 3; + bcast[1] = 1; + const TensorMap<Tensor<const float, 2>> input_tensor(input.data(), 3, 3); + Tensor<float, 2> output_tensor= (input_tensor - input_tensor.maximum(depth_dim).eval().reshape(dims2d).broadcast(bcast)); + + for (int i = 0; i < 3; ++i) { + for (int j = 0; j < 3; ++j) { + VERIFY_IS_APPROX(output(i, j), output_tensor(i, j)); + } + } +} + + void test_cxx11_tensor_forced_eval() { CALL_SUBTEST(test_simple()); + CALL_SUBTEST(test_const()); } diff --git a/unsupported/test/cxx11_tensor_image_patch.cpp b/unsupported/test/cxx11_tensor_image_patch.cpp index 55d35eac0..26854f5a4 100644 --- a/unsupported/test/cxx11_tensor_image_patch.cpp +++ b/unsupported/test/cxx11_tensor_image_patch.cpp @@ -28,6 +28,9 @@ static void test_simple_patch() VERIFY_IS_EQUAL(single_pixel_patch.dimension(4), 7); for (int i = 0; i < tensor.size(); ++i) { + if (tensor.data()[i] != single_pixel_patch.data()[i]) { + std::cout << "Mismatch detected at index " << i << " : " << tensor.data()[i] << " vs " << single_pixel_patch.data()[i] << std::endl; + } VERIFY_IS_EQUAL(single_pixel_patch.data()[i], tensor.data()[i]); } @@ -51,6 +54,9 @@ static void test_simple_patch() if (r-1+i >= 0 && c-2+j >= 0 && r-1+i < 3 && c-2+j < 5) { expected = tensor(d, r-1+i, c-2+j, b); } + if (entire_image_patch(d, r, c, patchId, b) != expected) { + std::cout << "Mismatch detected at index i=" << i << " j=" << j << " r=" << r << " c=" << c << " d=" << d << " b=" << b << std::endl; + } VERIFY_IS_EQUAL(entire_image_patch(d, r, c, patchId, b), expected); } } @@ -68,6 +74,11 @@ static void test_simple_patch() VERIFY_IS_EQUAL(twod_patch.dimension(3), 3*5); VERIFY_IS_EQUAL(twod_patch.dimension(4), 7); + // Based on the calculation described in TensorTraits.h, padding happens to be 0. + int row_padding = 0; + int col_padding = 0; + int stride = 1; + for (int i = 0; i < 3; ++i) { for (int j = 0; j < 5; ++j) { int patchId = i+3*j; @@ -76,8 +87,13 @@ static void test_simple_patch() for (int d = 0; d < 2; ++d) { for (int b = 0; b < 7; ++b) { float expected = 0.0f; - if (r-1+i >= 0 && c-1+j >= 0 && r-1+i < 3 && c-1+j < 5) { - expected = tensor(d, r-1+i, c-1+j, b); + int row_offset = r*stride + i - row_padding; + int col_offset = c*stride + j - col_padding; + if (row_offset >= 0 && col_offset >= 0 && row_offset < tensor.dimension(1) && col_offset < tensor.dimension(2)) { + expected = tensor(d, row_offset, col_offset, b); + } + if (twod_patch(d, r, c, patchId, b) != expected) { + std::cout << "Mismatch detected at index i=" << i << " j=" << j << " r=" << r << " c=" << c << " d=" << d << " b=" << b << std::endl; } VERIFY_IS_EQUAL(twod_patch(d, r, c, patchId, b), expected); } @@ -88,6 +104,156 @@ static void test_simple_patch() } } +// Verifies VALID padding (no padding) with incrementing values. +static void test_patch_padding_valid() +{ + int input_depth = 3; + int input_rows = 3; + int input_cols = 3; + int input_batches = 1; + int ksize = 2; // Corresponds to the Rows and Cols for tensor.extract_image_patches<>. + int stride = 2; // Only same stride is supported. + Tensor<float, 4> tensor(input_depth, input_rows, input_cols, input_batches); + // Initializes tensor with incrementing numbers. + for (int i = 0; i < tensor.size(); ++i) { + tensor.data()[i] = i + 1; + } + Tensor<float, 5> result = tensor.extract_image_patches(ksize, ksize, stride, stride, PADDING_VALID); + + VERIFY_IS_EQUAL(result.dimension(0), input_depth); // depth + VERIFY_IS_EQUAL(result.dimension(1), ksize); // kernel rows + VERIFY_IS_EQUAL(result.dimension(2), ksize); // kernel cols + VERIFY_IS_EQUAL(result.dimension(3), 1); // number of patches + VERIFY_IS_EQUAL(result.dimension(4), input_batches); // number of batches + + // No padding is carried out. + int row_padding = 0; + int col_padding = 0; + + for (int i = 0; (i+stride+ksize-1) < input_rows; i += stride) { // input rows + for (int j = 0; (j+stride+ksize-1) < input_cols; j += stride) { // input cols + int patchId = i+input_rows*j; + for (int r = 0; r < ksize; ++r) { // patch rows + for (int c = 0; c < ksize; ++c) { // patch cols + for (int d = 0; d < input_depth; ++d) { // depth + for (int b = 0; b < input_batches; ++b) { // batch + float expected = 0.0f; + int row_offset = r + i - row_padding; + int col_offset = c + j - col_padding; + if (row_offset >= 0 && col_offset >= 0 && row_offset < input_rows && col_offset < input_cols) { + expected = tensor(d, row_offset, col_offset, b); + } + if (result(d, r, c, patchId, b) != expected) { + std::cout << "Mismatch detected at index i=" << i << " j=" << j << " r=" << r << " c=" << c << " d=" << d << " b=" << b << std::endl; + } + VERIFY_IS_EQUAL(result(d, r, c, patchId, b), expected); + } + } + } + } + } + } +} + +// Verifies VALID padding (no padding) with the same value. +static void test_patch_padding_valid_same_value() +{ + int input_depth = 1; + int input_rows = 5; + int input_cols = 5; + int input_batches = 2; + int ksize = 3; // Corresponds to the Rows and Cols for tensor.extract_image_patches<>. + int stride = 2; // Only same stride is supported. + Tensor<float, 4> tensor(input_depth, input_rows, input_cols, input_batches); + tensor = tensor.constant(11.0f); + Tensor<float, 5> result = tensor.extract_image_patches(ksize, ksize, stride, stride, PADDING_VALID); + + VERIFY_IS_EQUAL(result.dimension(0), input_depth); // depth + VERIFY_IS_EQUAL(result.dimension(1), ksize); // kernel rows + VERIFY_IS_EQUAL(result.dimension(2), ksize); // kernel cols + VERIFY_IS_EQUAL(result.dimension(3), 4); // number of patches + VERIFY_IS_EQUAL(result.dimension(4), input_batches); // number of batches + + // No padding is carried out. + int row_padding = 0; + int col_padding = 0; + + for (int i = 0; (i+stride+ksize-1) <= input_rows; i += stride) { // input rows + for (int j = 0; (j+stride+ksize-1) <= input_cols; j += stride) { // input cols + int patchId = i+input_rows*j; + for (int r = 0; r < ksize; ++r) { // patch rows + for (int c = 0; c < ksize; ++c) { // patch cols + for (int d = 0; d < input_depth; ++d) { // depth + for (int b = 0; b < input_batches; ++b) { // batch + float expected = 0.0f; + int row_offset = r + i - row_padding; + int col_offset = c + j - col_padding; + if (row_offset >= 0 && col_offset >= 0 && row_offset < input_rows && col_offset < input_cols) { + expected = tensor(d, row_offset, col_offset, b); + } + if (result(d, r, c, patchId, b) != expected) { + std::cout << "Mismatch detected at index i=" << i << " j=" << j << " r=" << r << " c=" << c << " d=" << d << " b=" << b << std::endl; + } + VERIFY_IS_EQUAL(result(d, r, c, patchId, b), expected); + } + } + } + } + } + } +} + +// Verifies SAME padding. +static void test_patch_padding_same() +{ + int input_depth = 3; + int input_rows = 4; + int input_cols = 2; + int input_batches = 1; + int ksize = 2; // Corresponds to the Rows and Cols for tensor.extract_image_patches<>. + int stride = 2; // Only same stride is supported. + Tensor<float, 4> tensor(input_depth, input_rows, input_cols, input_batches); + // Initializes tensor with incrementing numbers. + for (int i = 0; i < tensor.size(); ++i) { + tensor.data()[i] = i + 1; + } + Tensor<float, 5> result = tensor.extract_image_patches(ksize, ksize, stride, stride, PADDING_SAME); + + VERIFY_IS_EQUAL(result.dimension(0), input_depth); // depth + VERIFY_IS_EQUAL(result.dimension(1), ksize); // kernel rows + VERIFY_IS_EQUAL(result.dimension(2), ksize); // kernel cols + VERIFY_IS_EQUAL(result.dimension(3), 2); // number of patches + VERIFY_IS_EQUAL(result.dimension(4), input_batches); // number of batches + + // Based on the calculation described in TensorTraits.h, padding happens to be + // 0. + int row_padding = 0; + int col_padding = 0; + + for (int i = 0; (i+stride+ksize-1) <= input_rows; i += stride) { // input rows + for (int j = 0; (j+stride+ksize-1) <= input_cols; j += stride) { // input cols + int patchId = i+input_rows*j; + for (int r = 0; r < ksize; ++r) { // patch rows + for (int c = 0; c < ksize; ++c) { // patch cols + for (int d = 0; d < input_depth; ++d) { // depth + for (int b = 0; b < input_batches; ++b) { // batch + float expected = 0.0f; + int row_offset = r*stride + i - row_padding; + int col_offset = c*stride + j - col_padding; + if (row_offset >= 0 && col_offset >= 0 && row_offset < input_rows && col_offset < input_cols) { + expected = tensor(d, row_offset, col_offset, b); + } + if (result(d, r, c, patchId, b) != expected) { + std::cout << "Mismatch detected at index i=" << i << " j=" << j << " r=" << r << " c=" << c << " d=" << d << " b=" << b << std::endl; + } + VERIFY_IS_EQUAL(result(d, r, c, patchId, b), expected); + } + } + } + } + } + } +} static void test_patch_no_extra_dim() { @@ -103,6 +269,9 @@ static void test_patch_no_extra_dim() VERIFY_IS_EQUAL(single_pixel_patch.dimension(3), 3*5); for (int i = 0; i < tensor.size(); ++i) { + if (tensor.data()[i] != single_pixel_patch.data()[i]) { + std::cout << "Mismatch detected at index " << i << " : " << tensor.data()[i] << " vs " << single_pixel_patch.data()[i] << std::endl; + } VERIFY_IS_EQUAL(single_pixel_patch.data()[i], tensor.data()[i]); } @@ -124,6 +293,9 @@ static void test_patch_no_extra_dim() if (r-1+i >= 0 && c-2+j >= 0 && r-1+i < 3 && c-2+j < 5) { expected = tensor(d, r-1+i, c-2+j); } + if (entire_image_patch(d, r, c, patchId) != expected) { + std::cout << "Mismatch detected at index i=" << i << " j=" << j << " r=" << r << " c=" << c << " d=" << d << std::endl; + } VERIFY_IS_EQUAL(entire_image_patch(d, r, c, patchId), expected); } } @@ -139,6 +311,11 @@ static void test_patch_no_extra_dim() VERIFY_IS_EQUAL(twod_patch.dimension(2), 2); VERIFY_IS_EQUAL(twod_patch.dimension(3), 3*5); + // Based on the calculation described in TensorTraits.h, padding happens to be 0. + int row_padding = 0; + int col_padding = 0; + int stride = 1; + for (int i = 0; i < 3; ++i) { for (int j = 0; j < 5; ++j) { int patchId = i+3*j; @@ -146,8 +323,13 @@ static void test_patch_no_extra_dim() for (int c = 0; c < 2; ++c) { for (int d = 0; d < 2; ++d) { float expected = 0.0f; - if (r-1+i >= 0 && c-1+j >= 0 && r-1+i < 3 && c-1+j < 5) { - expected = tensor(d, r-1+i, c-1+j); + int row_offset = r*stride + i - row_padding; + int col_offset = c*stride + j - col_padding; + if (row_offset >= 0 && col_offset >= 0 && row_offset < tensor.dimension(1) && col_offset < tensor.dimension(2)) { + expected = tensor(d, row_offset, col_offset); + } + if (twod_patch(d, r, c, patchId) != expected) { + std::cout << "Mismatch detected at index i=" << i << " j=" << j << " r=" << r << " c=" << c << " d=" << d << std::endl; } VERIFY_IS_EQUAL(twod_patch(d, r, c, patchId), expected); } @@ -181,6 +363,9 @@ static void test_imagenet_patches() if (r-5+i >= 0 && c-5+j >= 0 && r-5+i < 128 && c-5+j < 128) { expected = l_in(d, r-5+i, c-5+j, b); } + if (l_out(d, r, c, patchId, b) != expected) { + std::cout << "Mismatch detected at index i=" << i << " j=" << j << " r=" << r << " c=" << c << " d=" << d << " b=" << b << std::endl; + } VERIFY_IS_EQUAL(l_out(d, r, c, patchId, b), expected); } } @@ -208,6 +393,9 @@ static void test_imagenet_patches() if (r-4+i >= 0 && c-4+j >= 0 && r-4+i < 64 && c-4+j < 64) { expected = l_in(d, r-4+i, c-4+j, b); } + if (l_out(d, r, c, patchId, b) != expected) { + std::cout << "Mismatch detected at index i=" << i << " j=" << j << " r=" << r << " c=" << c << " d=" << d << " b=" << b << std::endl; + } VERIFY_IS_EQUAL(l_out(d, r, c, patchId, b), expected); } } @@ -235,6 +423,9 @@ static void test_imagenet_patches() if (r-3+i >= 0 && c-3+j >= 0 && r-3+i < 16 && c-3+j < 16) { expected = l_in(d, r-3+i, c-3+j, b); } + if (l_out(d, r, c, patchId, b) != expected) { + std::cout << "Mismatch detected at index i=" << i << " j=" << j << " r=" << r << " c=" << c << " d=" << d << " b=" << b << std::endl; + } VERIFY_IS_EQUAL(l_out(d, r, c, patchId, b), expected); } } @@ -262,6 +453,9 @@ static void test_imagenet_patches() if (r-1+i >= 0 && c-1+j >= 0 && r-1+i < 13 && c-1+j < 13) { expected = l_in(d, r-1+i, c-1+j, b); } + if (l_out(d, r, c, patchId, b) != expected) { + std::cout << "Mismatch detected at index i=" << i << " j=" << j << " r=" << r << " c=" << c << " d=" << d << " b=" << b << std::endl; + } VERIFY_IS_EQUAL(l_out(d, r, c, patchId, b), expected); } } @@ -271,10 +465,12 @@ static void test_imagenet_patches() } } - void test_cxx11_tensor_image_patch() { CALL_SUBTEST(test_simple_patch()); CALL_SUBTEST(test_patch_no_extra_dim()); + CALL_SUBTEST(test_patch_padding_valid()); + CALL_SUBTEST(test_patch_padding_valid_same_value()); + CALL_SUBTEST(test_patch_padding_same()); CALL_SUBTEST(test_imagenet_patches()); } diff --git a/unsupported/test/cxx11_tensor_map.cpp b/unsupported/test/cxx11_tensor_map.cpp index 478c20306..9cf2eb150 100644 --- a/unsupported/test/cxx11_tensor_map.cpp +++ b/unsupported/test/cxx11_tensor_map.cpp @@ -29,6 +29,7 @@ static void test_1d() vec1(4) = 23; vec2(4) = 4; vec1(5) = 42; vec2(5) = 5; + VERIFY_IS_EQUAL(vec1.rank(), 1); VERIFY_IS_EQUAL(vec1.size(), 6); VERIFY_IS_EQUAL(vec1.dimension(0), 6); @@ -69,10 +70,12 @@ static void test_2d() TensorMap<Tensor<const int, 2>> mat3(mat1.data(), 2, 3); TensorMap<Tensor<const int, 2, RowMajor>> mat4(mat2.data(), 2, 3); + VERIFY_IS_EQUAL(mat3.rank(), 2); VERIFY_IS_EQUAL(mat3.size(), 6); VERIFY_IS_EQUAL(mat3.dimension(0), 2); VERIFY_IS_EQUAL(mat3.dimension(1), 3); + VERIFY_IS_EQUAL(mat4.rank(), 2); VERIFY_IS_EQUAL(mat4.size(), 6); VERIFY_IS_EQUAL(mat4.dimension(0), 2); VERIFY_IS_EQUAL(mat4.dimension(1), 3); @@ -109,13 +112,15 @@ static void test_3d() } TensorMap<Tensor<const int, 3>> mat3(mat1.data(), 2, 3, 7); - TensorMap<Tensor<const int, 3, RowMajor>> mat4(mat2.data(), 2, 3, 7); + TensorMap<Tensor<const int, 3, RowMajor>> mat4(mat2.data(), array<DenseIndex, 3>{{2, 3, 7}}); + VERIFY_IS_EQUAL(mat3.rank(), 3); VERIFY_IS_EQUAL(mat3.size(), 2*3*7); VERIFY_IS_EQUAL(mat3.dimension(0), 2); VERIFY_IS_EQUAL(mat3.dimension(1), 3); VERIFY_IS_EQUAL(mat3.dimension(2), 7); + VERIFY_IS_EQUAL(mat4.rank(), 3); VERIFY_IS_EQUAL(mat4.size(), 2*3*7); VERIFY_IS_EQUAL(mat4.dimension(0), 2); VERIFY_IS_EQUAL(mat4.dimension(1), 3); diff --git a/unsupported/test/cxx11_tensor_morphing.cpp b/unsupported/test/cxx11_tensor_morphing.cpp index 78b0dade0..b4b0a55b6 100644 --- a/unsupported/test/cxx11_tensor_morphing.cpp +++ b/unsupported/test/cxx11_tensor_morphing.cpp @@ -89,19 +89,19 @@ static void test_reshape_as_lvalue() } } - +template<int DataLayout> static void test_simple_slice() { - Tensor<float, 5> tensor(2,3,5,7,11); + Tensor<float, 5, DataLayout> tensor(2,3,5,7,11); tensor.setRandom(); - Tensor<float, 5> slice1(1,1,1,1,1); + Tensor<float, 5, DataLayout> slice1(1,1,1,1,1); Eigen::DSizes<ptrdiff_t, 5> indices(1,2,3,4,5); Eigen::DSizes<ptrdiff_t, 5> sizes(1,1,1,1,1); slice1 = tensor.slice(indices, sizes); VERIFY_IS_EQUAL(slice1(0,0,0,0,0), tensor(1,2,3,4,5)); - Tensor<float, 5> slice2(1,1,2,2,3); + Tensor<float, 5, DataLayout> slice2(1,1,2,2,3); Eigen::DSizes<ptrdiff_t, 5> indices2(1,1,3,4,5); Eigen::DSizes<ptrdiff_t, 5> sizes2(1,1,2,2,3); slice2 = tensor.slice(indices2, sizes2); @@ -114,7 +114,7 @@ static void test_simple_slice() } } - +// TODO(andydavis) Add RowMajor support when TensorContract supports RowMajor. static void test_slice_in_expr() { MatrixXf m1(7,7); MatrixXf m2(3,3); @@ -141,21 +141,28 @@ static void test_slice_in_expr() { VERIFY_IS_APPROX(res(i,j), m3(i,j)); } } -} + // Take an arbitrary slice of an arbitrarily sized tensor. + TensorMap<Tensor<const float, 2>> tensor4(m1.data(), 7, 7); + Tensor<float, 1> tensor6 = tensor4.reshape(DSizes<ptrdiff_t, 1>(7*7)).exp().slice(DSizes<ptrdiff_t, 1>(0), DSizes<ptrdiff_t, 1>(35)); + for (int i = 0; i < 35; ++i) { + VERIFY_IS_APPROX(tensor6(i), expf(tensor4.data()[i])); + } +} +template<int DataLayout> static void test_slice_as_lvalue() { - Tensor<float, 3> tensor1(2,2,7); + Tensor<float, 3, DataLayout> tensor1(2,2,7); tensor1.setRandom(); - Tensor<float, 3> tensor2(2,2,7); + Tensor<float, 3, DataLayout> tensor2(2,2,7); tensor2.setRandom(); - Tensor<float, 3> tensor3(4,3,5); + Tensor<float, 3, DataLayout> tensor3(4,3,5); tensor3.setRandom(); - Tensor<float, 3> tensor4(4,3,2); + Tensor<float, 3, DataLayout> tensor4(4,3,2); tensor4.setRandom(); - Tensor<float, 3> result(4,5,7); + Tensor<float, 3, DataLayout> result(4,5,7); Eigen::DSizes<ptrdiff_t, 3> sizes12(2,2,7); Eigen::DSizes<ptrdiff_t, 3> first_slice(0,0,0); result.slice(first_slice, sizes12) = tensor1; @@ -190,10 +197,10 @@ static void test_slice_as_lvalue() } } - +template<int DataLayout> static void test_slice_raw_data() { - Tensor<float, 4> tensor(3,5,7,11); + Tensor<float, 4, DataLayout> tensor(3,5,7,11); tensor.setRandom(); Eigen::DSizes<ptrdiff_t, 4> offsets(1,2,3,4); @@ -203,40 +210,78 @@ static void test_slice_raw_data() VERIFY_IS_EQUAL(slice1.dimensions().TotalSize(), 1ul); VERIFY_IS_EQUAL(slice1.data()[0], tensor(1,2,3,4)); - extents = Eigen::DSizes<ptrdiff_t, 4>(2,1,1,1); - auto slice2 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice()); - VERIFY_IS_EQUAL(slice2.dimensions().TotalSize(), 2ul); - VERIFY_IS_EQUAL(slice2.data()[0], tensor(1,2,3,4)); - VERIFY_IS_EQUAL(slice2.data()[1], tensor(2,2,3,4)); + if (DataLayout == ColMajor) { + extents = Eigen::DSizes<ptrdiff_t, 4>(2,1,1,1); + auto slice2 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice()); + VERIFY_IS_EQUAL(slice2.dimensions().TotalSize(), 2ul); + VERIFY_IS_EQUAL(slice2.data()[0], tensor(1,2,3,4)); + VERIFY_IS_EQUAL(slice2.data()[1], tensor(2,2,3,4)); + } else { + extents = Eigen::DSizes<ptrdiff_t, 4>(1,1,1,2); + auto slice2 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice()); + VERIFY_IS_EQUAL(slice2.dimensions().TotalSize(), 2ul); + VERIFY_IS_EQUAL(slice2.data()[0], tensor(1,2,3,4)); + VERIFY_IS_EQUAL(slice2.data()[1], tensor(1,2,3,5)); + } extents = Eigen::DSizes<ptrdiff_t, 4>(1,2,1,1); auto slice3 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice()); VERIFY_IS_EQUAL(slice3.dimensions().TotalSize(), 2ul); VERIFY_IS_EQUAL(slice3.data(), static_cast<float*>(0)); - offsets = Eigen::DSizes<ptrdiff_t, 4>(0,2,3,4); - extents = Eigen::DSizes<ptrdiff_t, 4>(3,2,1,1); - auto slice4 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice()); - VERIFY_IS_EQUAL(slice4.dimensions().TotalSize(), 6ul); - for (int i = 0; i < 3; ++i) { - for (int j = 0; j < 2; ++j) { - VERIFY_IS_EQUAL(slice4.data()[i+3*j], tensor(i,2+j,3,4)); + if (DataLayout == ColMajor) { + offsets = Eigen::DSizes<ptrdiff_t, 4>(0,2,3,4); + extents = Eigen::DSizes<ptrdiff_t, 4>(3,2,1,1); + auto slice4 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice()); + VERIFY_IS_EQUAL(slice4.dimensions().TotalSize(), 6ul); + for (int i = 0; i < 3; ++i) { + for (int j = 0; j < 2; ++j) { + VERIFY_IS_EQUAL(slice4.data()[i+3*j], tensor(i,2+j,3,4)); + } + } + } else { + offsets = Eigen::DSizes<ptrdiff_t, 4>(1,2,3,0); + extents = Eigen::DSizes<ptrdiff_t, 4>(1,1,2,11); + auto slice4 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice()); + VERIFY_IS_EQUAL(slice4.dimensions().TotalSize(), 22ul); + for (int l = 0; l < 11; ++l) { + for (int k = 0; k < 2; ++k) { + VERIFY_IS_EQUAL(slice4.data()[l+11*k], tensor(1,2,3+k,l)); + } } } - offsets = Eigen::DSizes<ptrdiff_t, 4>(0,0,0,4); - extents = Eigen::DSizes<ptrdiff_t, 4>(3,5,7,2); - auto slice5 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice()); - VERIFY_IS_EQUAL(slice5.dimensions().TotalSize(), 210ul); - for (int i = 0; i < 3; ++i) { - for (int j = 0; j < 5; ++j) { + if (DataLayout == ColMajor) { + offsets = Eigen::DSizes<ptrdiff_t, 4>(0,0,0,4); + extents = Eigen::DSizes<ptrdiff_t, 4>(3,5,7,2); + auto slice5 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice()); + VERIFY_IS_EQUAL(slice5.dimensions().TotalSize(), 210ul); + for (int i = 0; i < 3; ++i) { + for (int j = 0; j < 5; ++j) { + for (int k = 0; k < 7; ++k) { + for (int l = 0; l < 2; ++l) { + int slice_index = i + 3 * (j + 5 * (k + 7 * l)); + VERIFY_IS_EQUAL(slice5.data()[slice_index], tensor(i,j,k,l+4)); + } + } + } + } + } else { + offsets = Eigen::DSizes<ptrdiff_t, 4>(1,0,0,0); + extents = Eigen::DSizes<ptrdiff_t, 4>(2,5,7,11); + auto slice5 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice()); + VERIFY_IS_EQUAL(slice5.dimensions().TotalSize(), 770ul); + for (int l = 0; l < 11; ++l) { for (int k = 0; k < 7; ++k) { - for (int l = 0; l < 2; ++l) { - int slice_index = i + 3 * (j + 5 * (k + 7 * l)); - VERIFY_IS_EQUAL(slice5.data()[slice_index], tensor(i,j,k,l+4)); + for (int j = 0; j < 5; ++j) { + for (int i = 0; i < 2; ++i) { + int slice_index = l + 11 * (k + 7 * (j + 5 * i)); + VERIFY_IS_EQUAL(slice5.data()[slice_index], tensor(i+1,j,k,l)); + } } } } + } offsets = Eigen::DSizes<ptrdiff_t, 4>(0,0,0,0); @@ -247,14 +292,38 @@ static void test_slice_raw_data() } +static void test_composition() +{ + Eigen::Tensor<float, 2> matrix(7, 11); + matrix.setRandom(); + + const DSizes<ptrdiff_t, 3> newDims{{1, 1, 11}}; + Eigen::Tensor<float, 3> tensor = + matrix.slice(DSizes<ptrdiff_t, 2>(2, 0), DSizes<ptrdiff_t, 2>(1, 11)).reshape(newDims); + + VERIFY_IS_EQUAL(tensor.dimensions().TotalSize(), 11ul); + VERIFY_IS_EQUAL(tensor.dimension(0), 1); + VERIFY_IS_EQUAL(tensor.dimension(1), 1); + VERIFY_IS_EQUAL(tensor.dimension(2), 11); + for (int i = 0; i < 11; ++i) { + VERIFY_IS_EQUAL(tensor(0,0,i), matrix(2,i)); + } +} + + void test_cxx11_tensor_morphing() { CALL_SUBTEST(test_simple_reshape()); CALL_SUBTEST(test_reshape_in_expr()); CALL_SUBTEST(test_reshape_as_lvalue()); - CALL_SUBTEST(test_simple_slice()); + CALL_SUBTEST(test_simple_slice<ColMajor>()); + CALL_SUBTEST(test_simple_slice<RowMajor>()); CALL_SUBTEST(test_slice_in_expr()); - CALL_SUBTEST(test_slice_as_lvalue()); - CALL_SUBTEST(test_slice_raw_data()); + CALL_SUBTEST(test_slice_as_lvalue<ColMajor>()); + CALL_SUBTEST(test_slice_as_lvalue<RowMajor>()); + CALL_SUBTEST(test_slice_raw_data<ColMajor>()); + CALL_SUBTEST(test_slice_raw_data<RowMajor>()); + + CALL_SUBTEST(test_composition()); } diff --git a/unsupported/test/cxx11_tensor_of_strings.cpp b/unsupported/test/cxx11_tensor_of_strings.cpp index 0ffa341c4..8d05d154e 100644 --- a/unsupported/test/cxx11_tensor_of_strings.cpp +++ b/unsupported/test/cxx11_tensor_of_strings.cpp @@ -8,19 +8,18 @@ // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. #include "main.h" -#include <string> + #include <Eigen/CXX11/Tensor> -using std::string; using Eigen::Tensor; using Eigen::TensorMap; static void test_assign() { - string data1[6]; - TensorMap<Tensor<string, 2>> mat1(data1, 2, 3); - string data2[6]; - const TensorMap<Tensor<const string, 2>> mat2(data2, 2, 3); + std::string data1[6]; + TensorMap<Tensor<std::string, 2>> mat1(data1, 2, 3); + std::string data2[6]; + const TensorMap<Tensor<const std::string, 2>> mat2(data2, 2, 3); for (int i = 0; i < 6; ++i) { std::ostringstream s1; @@ -31,16 +30,16 @@ static void test_assign() data2[i] = s2.str(); } - Tensor<string, 2> rslt1; + Tensor<std::string, 2> rslt1; rslt1 = mat1; - Tensor<string, 2> rslt2; + Tensor<std::string, 2> rslt2; rslt2 = mat2; - Tensor<string, 2> rslt3 = mat1; - Tensor<string, 2> rslt4 = mat2; + Tensor<std::string, 2> rslt3 = mat1; + Tensor<std::string, 2> rslt4 = mat2; - Tensor<string, 2> rslt5(mat1); - Tensor<string, 2> rslt6(mat2); + Tensor<std::string, 2> rslt5(mat1); + Tensor<std::string, 2> rslt6(mat2); for (int i = 0; i < 2; ++i) { for (int j = 0; j < 3; ++j) { @@ -57,8 +56,8 @@ static void test_assign() static void test_concat() { - Tensor<string, 2> t1(2, 3); - Tensor<string, 2> t2(2, 3); + Tensor<std::string, 2> t1(2, 3); + Tensor<std::string, 2> t2(2, 3); for (int i = 0; i < 2; ++i) { for (int j = 0; j < 3; ++j) { @@ -71,7 +70,7 @@ static void test_concat() } } - Tensor<string, 2> result = t1.concatenate(t2, 1); + Tensor<std::string, 2> result = t1.concatenate(t2, 1); VERIFY_IS_EQUAL(result.dimension(0), 2); VERIFY_IS_EQUAL(result.dimension(1), 6); @@ -86,7 +85,7 @@ static void test_concat() static void test_slices() { - Tensor<string, 2> data(2, 6); + Tensor<std::string, 2> data(2, 6); for (int i = 0; i < 2; ++i) { for (int j = 0; j < 3; ++j) { std::ostringstream s1; @@ -99,8 +98,8 @@ static void test_slices() const Eigen::DSizes<ptrdiff_t, 2> first_half{{0, 0}}; const Eigen::DSizes<ptrdiff_t, 2> second_half{{0, 3}}; - Tensor<string, 2> t1 = data.slice(first_half, half_size); - Tensor<string, 2> t2 = data.slice(second_half, half_size); + Tensor<std::string, 2> t1 = data.slice(first_half, half_size); + Tensor<std::string, 2> t2 = data.slice(second_half, half_size); for (int i = 0; i < 2; ++i) { for (int j = 0; j < 3; ++j) { @@ -113,8 +112,8 @@ static void test_slices() static void test_additions() { - Tensor<string, 1> data1(3); - Tensor<string, 1> data2(3); + Tensor<std::string, 1> data1(3); + Tensor<std::string, 1> data2(3); for (int i = 0; i < 3; ++i) { data1(i) = "abc"; std::ostringstream s1; @@ -122,16 +121,26 @@ static void test_additions() data2(i) = s1.str(); } - Tensor<string, 1> sum = data1 + data2; + Tensor<std::string, 1> sum = data1 + data2; for (int i = 0; i < 3; ++i) { std::ostringstream concat; concat << "abc" << i; - string expected = concat.str(); + std::string expected = concat.str(); VERIFY_IS_EQUAL(sum(i), expected); } } +static void test_initialization() +{ + Tensor<std::string, 2> a(2, 3); + a.setConstant(std::string("foo")); + for (int i = 0; i < 2*3; ++i) { + VERIFY_IS_EQUAL(a(i), std::string("foo")); + } +} + + void test_cxx11_tensor_of_strings() { // Beware: none of this is likely to ever work on a GPU. @@ -139,4 +148,5 @@ void test_cxx11_tensor_of_strings() CALL_SUBTEST(test_concat()); CALL_SUBTEST(test_slices()); CALL_SUBTEST(test_additions()); + CALL_SUBTEST(test_initialization()); } diff --git a/unsupported/test/cxx11_tensor_padding.cpp b/unsupported/test/cxx11_tensor_padding.cpp index 6f74216dd..ffa19896e 100644 --- a/unsupported/test/cxx11_tensor_padding.cpp +++ b/unsupported/test/cxx11_tensor_padding.cpp @@ -13,9 +13,10 @@ using Eigen::Tensor; +template<int DataLayout> static void test_simple_padding() { - Tensor<float, 4> tensor(2,3,5,7); + Tensor<float, 4, DataLayout> tensor(2,3,5,7); tensor.setRandom(); array<std::pair<ptrdiff_t, ptrdiff_t>, 4> paddings; @@ -24,7 +25,7 @@ static void test_simple_padding() paddings[2] = std::make_pair(3, 4); paddings[3] = std::make_pair(0, 0); - Tensor<float, 4> padded; + Tensor<float, 4, DataLayout> padded; padded = tensor.pad(paddings); VERIFY_IS_EQUAL(padded.dimension(0), 2+0); @@ -47,9 +48,10 @@ static void test_simple_padding() } } +template<int DataLayout> static void test_padded_expr() { - Tensor<float, 4> tensor(2,3,5,7); + Tensor<float, 4, DataLayout> tensor(2,3,5,7); tensor.setRandom(); array<std::pair<ptrdiff_t, ptrdiff_t>, 4> paddings; @@ -62,17 +64,19 @@ static void test_padded_expr() reshape_dims[0] = 12; reshape_dims[1] = 84; - Tensor<float, 2> result; + Tensor<float, 2, DataLayout> result; result = tensor.pad(paddings).reshape(reshape_dims); for (int i = 0; i < 2; ++i) { for (int j = 0; j < 6; ++j) { for (int k = 0; k < 12; ++k) { for (int l = 0; l < 7; ++l) { + const float result_value = DataLayout == ColMajor ? + result(i+2*j,k+12*l) : result(j+6*i,l+7*k); if (j >= 2 && j < 5 && k >= 3 && k < 8) { - VERIFY_IS_EQUAL(result(i+2*j,k+12*l), tensor(i,j-2,k-3,l)); + VERIFY_IS_EQUAL(result_value, tensor(i,j-2,k-3,l)); } else { - VERIFY_IS_EQUAL(result(i+2*j,k+12*l), 0.0f); + VERIFY_IS_EQUAL(result_value, 0.0f); } } } @@ -80,9 +84,10 @@ static void test_padded_expr() } } - void test_cxx11_tensor_padding() { - CALL_SUBTEST(test_simple_padding()); - CALL_SUBTEST(test_padded_expr()); + CALL_SUBTEST(test_simple_padding<ColMajor>()); + CALL_SUBTEST(test_simple_padding<RowMajor>()); + CALL_SUBTEST(test_padded_expr<ColMajor>()); + CALL_SUBTEST(test_padded_expr<RowMajor>()); } diff --git a/unsupported/test/cxx11_tensor_patch.cpp b/unsupported/test/cxx11_tensor_patch.cpp index e2ba5bfd8..0ee7b46d4 100644 --- a/unsupported/test/cxx11_tensor_patch.cpp +++ b/unsupported/test/cxx11_tensor_patch.cpp @@ -36,6 +36,23 @@ static void test_simple_patch() VERIFY_IS_EQUAL(tensor.data()[i], no_patch.data()[i]); } + patch_dims[0] = 2; + patch_dims[1] = 3; + patch_dims[2] = 5; + patch_dims[3] = 7; + Tensor<float, 5> single_patch; + single_patch = tensor.extract_patches(patch_dims); + + VERIFY_IS_EQUAL(single_patch.dimension(0), 2); + VERIFY_IS_EQUAL(single_patch.dimension(1), 3); + VERIFY_IS_EQUAL(single_patch.dimension(2), 5); + VERIFY_IS_EQUAL(single_patch.dimension(3), 7); + VERIFY_IS_EQUAL(single_patch.dimension(4), 1); + + for (int i = 0; i < tensor.size(); ++i) { + VERIFY_IS_EQUAL(tensor.data()[i], single_patch.data()[i]); + } + patch_dims[0] = 1; patch_dims[1] = 2; patch_dims[2] = 2; diff --git a/unsupported/test/cxx11_tensor_reduction.cpp b/unsupported/test/cxx11_tensor_reduction.cpp index da9885166..99e19eba4 100644 --- a/unsupported/test/cxx11_tensor_reduction.cpp +++ b/unsupported/test/cxx11_tensor_reduction.cpp @@ -13,15 +13,15 @@ using Eigen::Tensor; -static void test_simple_reductions() -{ - Tensor<float, 4> tensor(2,3,5,7); +template <int DataLayout> +static void test_simple_reductions() { + Tensor<float, 4, DataLayout> tensor(2, 3, 5, 7); tensor.setRandom(); array<ptrdiff_t, 2> reduction_axis; reduction_axis[0] = 1; reduction_axis[1] = 3; - Tensor<float, 2> result = tensor.sum(reduction_axis); + Tensor<float, 2, DataLayout> result = tensor.sum(reduction_axis); VERIFY_IS_EQUAL(result.dimension(0), 2); VERIFY_IS_EQUAL(result.dimension(1), 5); for (int i = 0; i < 2; ++i) { @@ -36,6 +36,53 @@ static void test_simple_reductions() } } + { + Tensor<float, 1, DataLayout> sum1 = tensor.sum(); + VERIFY_IS_EQUAL(sum1.dimension(0), 1); + + array<ptrdiff_t, 4> reduction_axis; + reduction_axis[0] = 0; + reduction_axis[1] = 1; + reduction_axis[2] = 2; + reduction_axis[3] = 3; + Tensor<float, 1, DataLayout> sum2 = tensor.sum(reduction_axis); + VERIFY_IS_EQUAL(sum2.dimension(0), 1); + + VERIFY_IS_APPROX(sum1(0), sum2(0)); + } + + reduction_axis[0] = 0; + reduction_axis[1] = 2; + result = tensor.prod(reduction_axis); + VERIFY_IS_EQUAL(result.dimension(0), 3); + VERIFY_IS_EQUAL(result.dimension(1), 7); + for (int i = 0; i < 3; ++i) { + for (int j = 0; j < 7; ++j) { + float prod = 1.0f; + for (int k = 0; k < 2; ++k) { + for (int l = 0; l < 5; ++l) { + prod *= tensor(k, i, l, j); + } + } + VERIFY_IS_APPROX(result(i, j), prod); + } + } + + { + Tensor<float, 1, DataLayout> prod1 = tensor.prod(); + VERIFY_IS_EQUAL(prod1.dimension(0), 1); + + array<ptrdiff_t, 4> reduction_axis; + reduction_axis[0] = 0; + reduction_axis[1] = 1; + reduction_axis[2] = 2; + reduction_axis[3] = 3; + Tensor<float, 1, DataLayout> prod2 = tensor.prod(reduction_axis); + VERIFY_IS_EQUAL(prod2.dimension(0), 1); + + VERIFY_IS_APPROX(prod1(0), prod2(0)); + } + reduction_axis[0] = 0; reduction_axis[1] = 2; result = tensor.maximum(reduction_axis); @@ -53,6 +100,21 @@ static void test_simple_reductions() } } + { + Tensor<float, 1, DataLayout> max1 = tensor.maximum(); + VERIFY_IS_EQUAL(max1.dimension(0), 1); + + array<ptrdiff_t, 4> reduction_axis; + reduction_axis[0] = 0; + reduction_axis[1] = 1; + reduction_axis[2] = 2; + reduction_axis[3] = 3; + Tensor<float, 1, DataLayout> max2 = tensor.maximum(reduction_axis); + VERIFY_IS_EQUAL(max2.dimension(0), 1); + + VERIFY_IS_APPROX(max1(0), max2(0)); + } + reduction_axis[0] = 0; reduction_axis[1] = 1; result = tensor.minimum(reduction_axis); @@ -63,24 +125,72 @@ static void test_simple_reductions() float min_val = (std::numeric_limits<float>::max)(); for (int k = 0; k < 2; ++k) { for (int l = 0; l < 3; ++l) { - min_val = (std::min)(min_val, tensor(k, l, i, j)); + min_val = (std::min)(min_val, tensor(k, l, i, j)); } } VERIFY_IS_APPROX(result(i, j), min_val); } } -} + { + Tensor<float, 1, DataLayout> min1 = tensor.minimum(); + VERIFY_IS_EQUAL(min1.dimension(0), 1); + + array<ptrdiff_t, 4> reduction_axis; + reduction_axis[0] = 0; + reduction_axis[1] = 1; + reduction_axis[2] = 2; + reduction_axis[3] = 3; + Tensor<float, 1, DataLayout> min2 = tensor.minimum(reduction_axis); + VERIFY_IS_EQUAL(min2.dimension(0), 1); -static void test_full_reductions() -{ - Tensor<float, 2> tensor(2,3); + VERIFY_IS_APPROX(min1(0), min2(0)); + } + + reduction_axis[0] = 0; + reduction_axis[1] = 1; + result = tensor.mean(reduction_axis); + VERIFY_IS_EQUAL(result.dimension(0), 5); + VERIFY_IS_EQUAL(result.dimension(1), 7); + for (int i = 0; i < 5; ++i) { + for (int j = 0; j < 7; ++j) { + float sum = 0.0f; + int count = 0; + for (int k = 0; k < 2; ++k) { + for (int l = 0; l < 3; ++l) { + sum += tensor(k, l, i, j); + ++count; + } + } + VERIFY_IS_APPROX(result(i, j), sum / count); + } + } + + { + Tensor<float, 1, DataLayout> mean1 = tensor.mean(); + VERIFY_IS_EQUAL(mean1.dimension(0), 1); + + array<ptrdiff_t, 4> reduction_axis; + reduction_axis[0] = 0; + reduction_axis[1] = 1; + reduction_axis[2] = 2; + reduction_axis[3] = 3; + Tensor<float, 1, DataLayout> mean2 = tensor.mean(reduction_axis); + VERIFY_IS_EQUAL(mean2.dimension(0), 1); + + VERIFY_IS_APPROX(mean1(0), mean2(0)); + } +} + +template <int DataLayout> +static void test_full_reductions() { + Tensor<float, 2, DataLayout> tensor(2, 3); tensor.setRandom(); array<ptrdiff_t, 2> reduction_axis; reduction_axis[0] = 0; reduction_axis[1] = 1; - Tensor<float, 1> result = tensor.sum(reduction_axis); + Tensor<float, 1, DataLayout> result = tensor.sum(reduction_axis); VERIFY_IS_EQUAL(result.dimension(0), 1); float sum = 0.0f; @@ -103,30 +213,26 @@ static void test_full_reductions() VERIFY_IS_APPROX(result(0), sqrtf(sum)); } - struct UserReducer { - UserReducer(float offset) : offset_(offset), sum_(0.0f) {} - void reduce(const float val) { - sum_ += val * val; - } - float finalize() const { - return 1.0f / (sum_ + offset_); - } + static const bool PacketAccess = false; + UserReducer(float offset) : offset_(offset) {} + void reduce(const float val, float* accum) { *accum += val * val; } + float initialize() const { return 0; } + float finalize(const float accum) const { return 1.0f / (accum + offset_); } private: - float offset_; - float sum_; + const float offset_; }; -static void test_user_defined_reductions() -{ - Tensor<float, 2> tensor(5,7); +template <int DataLayout> +static void test_user_defined_reductions() { + Tensor<float, 2, DataLayout> tensor(5, 7); tensor.setRandom(); array<ptrdiff_t, 1> reduction_axis; reduction_axis[0] = 1; UserReducer reducer(10.0f); - Tensor<float, 1> result = tensor.reduce(reduction_axis, reducer); + Tensor<float, 1, DataLayout> result = tensor.reduce(reduction_axis, reducer); VERIFY_IS_EQUAL(result.dimension(0), 5); for (int i = 0; i < 5; ++i) { float expected = 10.0f; @@ -138,22 +244,24 @@ static void test_user_defined_reductions() } } - -static void test_tensor_maps() -{ - int inputs[2*3*5*7]; - TensorMap<Tensor<int, 4> > tensor_map(inputs, 2,3,5,7); - TensorMap<Tensor<const int, 4> > tensor_map_const(inputs, 2,3,5,7); - const TensorMap<Tensor<const int, 4> > tensor_map_const_const(inputs, 2,3,5,7); +template <int DataLayout> +static void test_tensor_maps() { + int inputs[2 * 3 * 5 * 7]; + TensorMap<Tensor<int, 4, DataLayout> > tensor_map(inputs, 2, 3, 5, 7); + TensorMap<Tensor<const int, 4, DataLayout> > tensor_map_const(inputs, 2, 3, 5, + 7); + const TensorMap<Tensor<const int, 4, DataLayout> > tensor_map_const_const( + inputs, 2, 3, 5, 7); tensor_map.setRandom(); array<ptrdiff_t, 2> reduction_axis; reduction_axis[0] = 1; reduction_axis[1] = 3; - Tensor<int, 2> result = tensor_map.sum(reduction_axis); - Tensor<int, 2> result2 = tensor_map_const.sum(reduction_axis); - Tensor<int, 2> result3 = tensor_map_const_const.sum(reduction_axis); + Tensor<int, 2, DataLayout> result = tensor_map.sum(reduction_axis); + Tensor<int, 2, DataLayout> result2 = tensor_map_const.sum(reduction_axis); + Tensor<int, 2, DataLayout> result3 = + tensor_map_const_const.sum(reduction_axis); for (int i = 0; i < 2; ++i) { for (int j = 0; j < 5; ++j) { @@ -170,11 +278,110 @@ static void test_tensor_maps() } } +template <int DataLayout> +static void test_static_dims() { + Tensor<float, 4, DataLayout> in(72, 53, 97, 113); + Tensor<float, 2, DataLayout> out(72, 97); + in.setRandom(); + +#if __cplusplus <= 199711L + array<int, 2> reduction_axis; + reduction_axis[0] = 1; + reduction_axis[1] = 3; +#else + Eigen::IndexList<Eigen::type2index<1>, Eigen::type2index<3> > reduction_axis; +#endif + + out = in.maximum(reduction_axis); + + for (int i = 0; i < 72; ++i) { + for (int j = 0; j < 97; ++j) { + float expected = -1e10f; + for (int k = 0; k < 53; ++k) { + for (int l = 0; l < 113; ++l) { + expected = (std::max)(expected, in(i, k, j, l)); + } + } + VERIFY_IS_APPROX(out(i, j), expected); + } + } +} + +template <int DataLayout> +static void test_innermost_last_dims() { + Tensor<float, 4, DataLayout> in(72, 53, 97, 113); + Tensor<float, 2, DataLayout> out(97, 113); + in.setRandom(); + +// Reduce on the innermost dimensions. +#if __cplusplus <= 199711L + array<int, 2> reduction_axis; + reduction_axis[0] = 0; + reduction_axis[1] = 1; +#else + // This triggers the use of packets for ColMajor. + Eigen::IndexList<Eigen::type2index<0>, Eigen::type2index<1> > reduction_axis; +#endif + + out = in.maximum(reduction_axis); + + for (int i = 0; i < 97; ++i) { + for (int j = 0; j < 113; ++j) { + float expected = -1e10f; + for (int k = 0; k < 53; ++k) { + for (int l = 0; l < 72; ++l) { + expected = (std::max)(expected, in(l, k, i, j)); + } + } + VERIFY_IS_APPROX(out(i, j), expected); + } + } +} + +template <int DataLayout> +static void test_innermost_first_dims() { + Tensor<float, 4, DataLayout> in(72, 53, 97, 113); + Tensor<float, 2, DataLayout> out(72, 53); + in.setRandom(); + +// Reduce on the innermost dimensions. +#if __cplusplus <= 199711L + array<int, 2> reduction_axis; + reduction_axis[0] = 2; + reduction_axis[1] = 3; +#else + // This triggers the use of packets for RowMajor. + Eigen::IndexList<Eigen::type2index<2>, Eigen::type2index<3>> reduction_axis; +#endif + + out = in.maximum(reduction_axis); + + for (int i = 0; i < 72; ++i) { + for (int j = 0; j < 53; ++j) { + float expected = -1e10f; + for (int k = 0; k < 97; ++k) { + for (int l = 0; l < 113; ++l) { + expected = (std::max)(expected, in(i, j, k, l)); + } + } + VERIFY_IS_APPROX(out(i, j), expected); + } + } +} -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()); +void test_cxx11_tensor_reduction() { + CALL_SUBTEST(test_simple_reductions<ColMajor>()); + CALL_SUBTEST(test_simple_reductions<RowMajor>()); + CALL_SUBTEST(test_full_reductions<ColMajor>()); + CALL_SUBTEST(test_full_reductions<RowMajor>()); + CALL_SUBTEST(test_user_defined_reductions<ColMajor>()); + CALL_SUBTEST(test_user_defined_reductions<RowMajor>()); + CALL_SUBTEST(test_tensor_maps<ColMajor>()); + CALL_SUBTEST(test_tensor_maps<RowMajor>()); + CALL_SUBTEST(test_static_dims<ColMajor>()); + CALL_SUBTEST(test_static_dims<RowMajor>()); + CALL_SUBTEST(test_innermost_last_dims<RowMajor>()); + CALL_SUBTEST(test_innermost_last_dims<ColMajor>()); + CALL_SUBTEST(test_innermost_first_dims<RowMajor>()); + CALL_SUBTEST(test_innermost_first_dims<ColMajor>()); } diff --git a/unsupported/test/cxx11_tensor_shuffling.cpp b/unsupported/test/cxx11_tensor_shuffling.cpp index 39c623499..ec623e1f9 100644 --- a/unsupported/test/cxx11_tensor_shuffling.cpp +++ b/unsupported/test/cxx11_tensor_shuffling.cpp @@ -14,9 +14,10 @@ using Eigen::Tensor; using Eigen::array; +template <int DataLayout> static void test_simple_shuffling() { - Tensor<float, 4> tensor(2,3,5,7); + Tensor<float, 4, DataLayout> tensor(2,3,5,7); tensor.setRandom(); array<ptrdiff_t, 4> shuffles; shuffles[0] = 0; @@ -24,7 +25,7 @@ static void test_simple_shuffling() shuffles[2] = 2; shuffles[3] = 3; - Tensor<float, 4> no_shuffle; + Tensor<float, 4, DataLayout> no_shuffle; no_shuffle = tensor.shuffle(shuffles); VERIFY_IS_EQUAL(no_shuffle.dimension(0), 2); @@ -46,7 +47,7 @@ static void test_simple_shuffling() shuffles[1] = 3; shuffles[2] = 1; shuffles[3] = 0; - Tensor<float, 4> shuffle; + Tensor<float, 4, DataLayout> shuffle; shuffle = tensor.shuffle(shuffles); VERIFY_IS_EQUAL(shuffle.dimension(0), 5); @@ -66,9 +67,10 @@ static void test_simple_shuffling() } +template <int DataLayout> static void test_expr_shuffling() { - Tensor<float, 4> tensor(2,3,5,7); + Tensor<float, 4, DataLayout> tensor(2,3,5,7); tensor.setRandom(); array<ptrdiff_t, 4> shuffles; @@ -76,10 +78,10 @@ static void test_expr_shuffling() shuffles[1] = 3; shuffles[2] = 1; shuffles[3] = 0; - Tensor<float, 4> expected; + Tensor<float, 4, DataLayout> expected; expected = tensor.shuffle(shuffles); - Tensor<float, 4> result(5,7,3,2); + Tensor<float, 4, DataLayout> result(5,7,3,2); array<int, 4> src_slice_dim{{2,3,1,7}}; array<int, 4> src_slice_start{{0,0,0,0}}; @@ -128,16 +130,17 @@ static void test_expr_shuffling() } +template <int DataLayout> static void test_shuffling_as_value() { - Tensor<float, 4> tensor(2,3,5,7); + Tensor<float, 4, DataLayout> tensor(2,3,5,7); tensor.setRandom(); array<ptrdiff_t, 4> shuffles; shuffles[2] = 0; shuffles[3] = 1; shuffles[1] = 2; shuffles[0] = 3; - Tensor<float, 4> shuffle(5,7,3,2); + Tensor<float, 4, DataLayout> shuffle(5,7,3,2); shuffle.shuffle(shuffles) = tensor; VERIFY_IS_EQUAL(shuffle.dimension(0), 5); @@ -158,7 +161,10 @@ static void test_shuffling_as_value() void test_cxx11_tensor_shuffling() { - CALL_SUBTEST(test_simple_shuffling()); - CALL_SUBTEST(test_expr_shuffling()); - CALL_SUBTEST(test_shuffling_as_value()); + CALL_SUBTEST(test_simple_shuffling<ColMajor>()); + CALL_SUBTEST(test_simple_shuffling<RowMajor>()); + CALL_SUBTEST(test_expr_shuffling<ColMajor>()); + CALL_SUBTEST(test_expr_shuffling<RowMajor>()); + CALL_SUBTEST(test_shuffling_as_value<ColMajor>()); + CALL_SUBTEST(test_shuffling_as_value<RowMajor>()); } diff --git a/unsupported/test/cxx11_tensor_simple.cpp b/unsupported/test/cxx11_tensor_simple.cpp index a70591c82..23855fca0 100644 --- a/unsupported/test/cxx11_tensor_simple.cpp +++ b/unsupported/test/cxx11_tensor_simple.cpp @@ -32,6 +32,7 @@ static void test_1d() vec1(5) = 42; vec2(5) = 5; vec3(5) = 0; vec4.setZero(); + VERIFY_IS_EQUAL((vec1.rank()), 1); VERIFY_IS_EQUAL((vec1.size()), 6); VERIFY_IS_EQUAL((vec1.dimensions()[0]), 6); @@ -99,10 +100,12 @@ static void test_2d() mat2(1,1) = 4; mat2(1,2) = 5; + VERIFY_IS_EQUAL((mat1.rank()), 2); VERIFY_IS_EQUAL((mat1.size()), 6); VERIFY_IS_EQUAL((mat1.dimensions()[0]), 2); VERIFY_IS_EQUAL((mat1.dimensions()[1]), 3); + VERIFY_IS_EQUAL((mat2.rank()), 2); VERIFY_IS_EQUAL((mat2.size()), 6); VERIFY_IS_EQUAL((mat2.dimensions()[0]), 2); VERIFY_IS_EQUAL((mat2.dimensions()[1]), 3); diff --git a/unsupported/test/cxx11_tensor_striding.cpp b/unsupported/test/cxx11_tensor_striding.cpp index 502569d1d..1feb39dca 100644 --- a/unsupported/test/cxx11_tensor_striding.cpp +++ b/unsupported/test/cxx11_tensor_striding.cpp @@ -13,9 +13,10 @@ using Eigen::Tensor; +template<int DataLayout> static void test_simple_striding() { - Tensor<float, 4> tensor(2,3,5,7); + Tensor<float, 4, DataLayout> tensor(2,3,5,7); tensor.setRandom(); array<ptrdiff_t, 4> strides; strides[0] = 1; @@ -23,7 +24,7 @@ static void test_simple_striding() strides[2] = 1; strides[3] = 1; - Tensor<float, 4> no_stride; + Tensor<float, 4, DataLayout> no_stride; no_stride = tensor.stride(strides); VERIFY_IS_EQUAL(no_stride.dimension(0), 2); @@ -45,7 +46,7 @@ static void test_simple_striding() strides[1] = 4; strides[2] = 2; strides[3] = 3; - Tensor<float, 4> stride; + Tensor<float, 4, DataLayout> stride; stride = tensor.stride(strides); VERIFY_IS_EQUAL(stride.dimension(0), 1); @@ -65,7 +66,36 @@ static void test_simple_striding() } +template<int DataLayout> +static void test_striding_as_lvalue() +{ + Tensor<float, 4, DataLayout> tensor(2,3,5,7); + tensor.setRandom(); + array<ptrdiff_t, 4> strides; + strides[0] = 2; + strides[1] = 4; + strides[2] = 2; + strides[3] = 3; + + Tensor<float, 4, DataLayout> result(3, 12, 10, 21); + result.stride(strides) = tensor; + + for (int i = 0; i < 2; ++i) { + for (int j = 0; j < 3; ++j) { + for (int k = 0; k < 5; ++k) { + for (int l = 0; l < 7; ++l) { + VERIFY_IS_EQUAL(tensor(i,j,k,l), result(2*i,4*j,2*k,3*l)); + } + } + } + } +} + + void test_cxx11_tensor_striding() { - CALL_SUBTEST(test_simple_striding()); + CALL_SUBTEST(test_simple_striding<ColMajor>()); + CALL_SUBTEST(test_simple_striding<RowMajor>()); + CALL_SUBTEST(test_striding_as_lvalue<ColMajor>()); + CALL_SUBTEST(test_striding_as_lvalue<RowMajor>()); } diff --git a/unsupported/test/cxx11_tensor_thread_pool.cpp b/unsupported/test/cxx11_tensor_thread_pool.cpp index f0de61f8b..e25912279 100644 --- a/unsupported/test/cxx11_tensor_thread_pool.cpp +++ b/unsupported/test/cxx11_tensor_thread_pool.cpp @@ -9,11 +9,11 @@ #define EIGEN_USE_THREADS -#include <iostream> + #include "main.h" +#include <iostream> #include <Eigen/CXX11/Tensor> - using Eigen::Tensor; static void test_multithread_elementwise() @@ -60,12 +60,12 @@ static void test_multithread_compound_assignment() } } - +template<int DataLayout> static void test_multithread_contraction() { - Tensor<float, 4> t_left(30, 50, 37, 31); - Tensor<float, 5> t_right(37, 31, 70, 2, 10); - Tensor<float, 5> t_result(30, 50, 70, 2, 10); + Tensor<float, 4, DataLayout> t_left(30, 50, 37, 31); + Tensor<float, 5, DataLayout> t_right(37, 31, 70, 2, 10); + Tensor<float, 5, DataLayout> t_result(30, 50, 70, 2, 10); t_left.setRandom(); t_right.setRandom(); @@ -74,11 +74,10 @@ static void test_multithread_contraction() typedef Tensor<float, 1>::DimensionPair DimPair; Eigen::array<DimPair, 2> dims({{DimPair(2, 0), DimPair(3, 1)}}); - - typedef Map<MatrixXf> MapXf; + typedef Map<Matrix<float, Dynamic, Dynamic, DataLayout>> MapXf; MapXf m_left(t_left.data(), 1500, 1147); MapXf m_right(t_right.data(), 1147, 1400); - MatrixXf m_result(1500, 1400); + Matrix<float, Dynamic, Dynamic, DataLayout> m_result(1500, 1400); Eigen::ThreadPoolDevice thread_pool_device(4); @@ -95,12 +94,12 @@ static void test_multithread_contraction() } } - +template<int DataLayout> static void test_contraction_corner_cases() { - Tensor<float, 2> t_left(32, 500); - Tensor<float, 2> t_right(32, 28*28); - Tensor<float, 2> t_result(500, 28*28); + Tensor<float, 2, DataLayout> t_left(32, 500); + Tensor<float, 2, DataLayout> t_right(32, 28*28); + Tensor<float, 2, DataLayout> t_result(500, 28*28); t_left = (t_left.constant(-0.5f) + t_left.random()) * 2.0f; t_right = (t_right.constant(-0.6f) + t_right.random()) * 2.0f; @@ -110,10 +109,10 @@ static void test_contraction_corner_cases() typedef Tensor<float, 1>::DimensionPair DimPair; Eigen::array<DimPair, 1> dims{{DimPair(0, 0)}}; - typedef Map<MatrixXf> MapXf; + typedef Map<Matrix<float, Dynamic, Dynamic, DataLayout>> MapXf; MapXf m_left(t_left.data(), 32, 500); MapXf m_right(t_right.data(), 32, 28*28); - MatrixXf m_result(500, 28*28); + Matrix<float, Dynamic, Dynamic, DataLayout> m_result(500, 28*28); Eigen::ThreadPoolDevice thread_pool_device(12); @@ -181,18 +180,18 @@ static void test_contraction_corner_cases() } } - +template<int DataLayout> static void test_multithread_contraction_agrees_with_singlethread() { int contract_size = internal::random<int>(1, 5000); - Tensor<float, 3> left(internal::random<int>(1, 80), - contract_size, - internal::random<int>(1, 100)); + Tensor<float, 3, DataLayout> left(internal::random<int>(1, 80), + contract_size, + internal::random<int>(1, 100)); - Tensor<float, 4> right(internal::random<int>(1, 25), - internal::random<int>(1, 37), - contract_size, - internal::random<int>(1, 51)); + Tensor<float, 4, DataLayout> right(internal::random<int>(1, 25), + internal::random<int>(1, 37), + contract_size, + internal::random<int>(1, 51)); left.setRandom(); right.setRandom(); @@ -206,13 +205,13 @@ static void test_multithread_contraction_agrees_with_singlethread() { Eigen::ThreadPoolDevice thread_pool_device(internal::random<int>(2, 11)); - Tensor<float, 5> st_result; + Tensor<float, 5, DataLayout> st_result; st_result = left.contract(right, dims); - Tensor<float, 5> tp_result(st_result.dimensions()); + Tensor<float, 5, DataLayout> tp_result(st_result.dimensions()); tp_result.device(thread_pool_device) = left.contract(right, dims); - VERIFY(internal::dimensions_match(st_result.dimensions(), tp_result.dimensions())); + VERIFY(dimensions_match(st_result.dimensions(), tp_result.dimensions())); for (ptrdiff_t i = 0; i < st_result.size(); i++) { // if both of the values are very small, then do nothing (because the test will fail // due to numerical precision issues when values are small) @@ -241,17 +240,30 @@ static void test_memcpy() { } +static void test_multithread_random() +{ + Eigen::ThreadPoolDevice device(2); + Tensor<float, 1> t(1 << 20); + t.device(device) = t.random<Eigen::internal::NormalRandomGenerator<float>>(); +} + + void test_cxx11_tensor_thread_pool() { CALL_SUBTEST(test_multithread_elementwise()); CALL_SUBTEST(test_multithread_compound_assignment()); - CALL_SUBTEST(test_multithread_contraction()); + CALL_SUBTEST(test_multithread_contraction<ColMajor>()); + CALL_SUBTEST(test_multithread_contraction<RowMajor>()); - CALL_SUBTEST(test_multithread_contraction_agrees_with_singlethread()); + CALL_SUBTEST(test_multithread_contraction_agrees_with_singlethread<ColMajor>()); + CALL_SUBTEST(test_multithread_contraction_agrees_with_singlethread<RowMajor>()); // Exercise various cases that have been problematic in the past. - CALL_SUBTEST(test_contraction_corner_cases()); + CALL_SUBTEST(test_contraction_corner_cases<ColMajor>()); + CALL_SUBTEST(test_contraction_corner_cases<RowMajor>()); CALL_SUBTEST(test_memcpy()); + + CALL_SUBTEST(test_multithread_random()); } |