Created many additional tests

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()

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());
 }

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_simple_broadcasting<ColMajor>());
-   CALL_SUBTEST(test_vectorized_broadcasting());
+  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>());
 }

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;
+  Tensor<float, 4, DataLayout> chip1;
-  chip1 = tensor.chip<0>(1);
+  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<float, 5, DataLayout> tensor = input1;
-  tensor.chip<0>(1) = input2;
+  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;
+  typedef TensorEvaluator<decltype(tensor.template chip<4>(3)), DefaultDevice> Evaluator4;
-  auto chip = Evaluator4(tensor.chip<4>(3), DefaultDevice());
+  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;
+  typedef TensorEvaluator<decltype(tensor.template chip<0>(0)), DefaultDevice> Evaluator0;
-  auto chip0 = Evaluator0(tensor.chip<0>(0), DefaultDevice());
+  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;
+  typedef TensorEvaluator<decltype(tensor.template chip<1>(0)), DefaultDevice> Evaluator1;
-  auto chip1 = Evaluator1(tensor.chip<1>(0), DefaultDevice());
+  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;
+  typedef TensorEvaluator<decltype(tensor.template chip<2>(0)), DefaultDevice> Evaluator2;
-  auto chip2 = Evaluator2(tensor.chip<2>(0), DefaultDevice());
+  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;
+  typedef TensorEvaluator<decltype(tensor.template chip<3>(0)), DefaultDevice> Evaluator3;
-  auto chip3 = Evaluator3(tensor.chip<3>(0), DefaultDevice());
+  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_simple_chip<ColMajor>());
-  CALL_SUBTEST(test_chip_in_expr());
+  CALL_SUBTEST(test_simple_chip<RowMajor>());
-  CALL_SUBTEST(test_chip_as_lvalue());
+  CALL_SUBTEST(test_dynamic_chip<ColMajor>());
-  CALL_SUBTEST(test_chip_raw_data());
+  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>());
 }

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, DataLayout> left(2, 3, 1);
-  Tensor<int, 3> right(3, 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, 2, DataLayout> left(2, 3);
-  Tensor<int, 3> right(2, 3, 1);
+  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, DataLayout> left(2, 3, 1);
-  Tensor<int, 3> right(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_dimension_failures<ColMajor>());
-   CALL_SUBTEST(test_static_dimension_failure());
+   CALL_SUBTEST(test_dimension_failures<RowMajor>());
-   CALL_SUBTEST(test_simple_concatenation());
+   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());
 }

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]));
+      }
+}

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, DataLayout> mat1(2, 3);
-  Tensor<float, 2> mat2(2, 3);
+  Tensor<float, 2, DataLayout> mat2(2, 3);
-  Tensor<float, 2> mat3(3, 2);
+  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, DataLayout> vec1({6});
-  Tensor<float, 1> vec2({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, 3, DataLayout> mat1(2, 2, 2);
-  Tensor<float, 4> mat2(2, 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, 4, DataLayout> t1(2, 5, 7, 3);
-  Tensor<float, 5> t2(2, 7, 11, 13, 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, 2, DataLayout> t1(2, 2);
-  Tensor<float, 3> t2(2, 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, DataLayout> t1(2, 2);
-  Tensor<float, 2> t2(2, 2);
+  Tensor<float, 2, DataLayout> t2(2, 2);
-  Tensor<float, 2> t3(2, 2);
+  Tensor<float, 2, DataLayout> t3(2, 2);
-  Tensor<float, 2> t4(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<Eigen::Matrix<float, Dynamic, Dynamic, DataLayout>>
-  Eigen::Map<MatrixXf> m2(t2.data(), 2, 2);
+      m1(t1.data(), 2, 2), m2(t2.data(), 2, 2), m3(t3.data(), 2, 2),
-  Eigen::Map<MatrixXf> m3(t3.data(), 2, 2);
+      m4(t4.data(), 2, 2);
-  Eigen::Map<MatrixXf> m4(t4.data(), 2, 2);
+  Eigen::Matrix<float, Dynamic, Dynamic, DataLayout>
-  Eigen::MatrixXf expected = (m1 * m4) * (m3 - m1 * m2);
+      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, DataLayout> mat1(2, 3);
-  Tensor<float, 2> mat2(3, 2);
+  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, DataLayout> mat1(2, 2, 2);
-  Tensor<float, 3> mat2(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, 3, DataLayout> mat1(4, 3, 5);
-  Tensor<float, 5> mat2(3, 2, 1, 5, 4);
+  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, DataLayout> mat3(5, 2, 1, 5);
-  Tensor<float, 4> mat4(2, 1, 5, 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++) {
+  if (DataLayout == ColMajor) {
-    for (size_t j = 0; j < 10; j++) {
+    for (size_t i = 0; i < 5; i++) {
-      VERIFY_IS_APPROX(mat3.data()[i + 5 * j], mat4.data()[j + 10 * 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, 4, DataLayout> t_left(30, 50, 8, 31);
-  Tensor<float, 5> t_right(8, 31, 7, 20, 10);
+  Tensor<float, 5, DataLayout> t_right(8, 31, 7, 20, 10);
-  Tensor<float, 5> t_result(30, 50, 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, 2, DataLayout> t_left(30, 50);
-  Tensor<float, 1> t_right(50);
+  Tensor<float, 1, DataLayout> t_right(50);
-  Tensor<float, 1> t_result(30);
+  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, 3, DataLayout> t_left(7, 13, 17);
-  Tensor<float, 2> t_right(1, 7);
+  Tensor<float, 2, DataLayout> t_right(1, 7);
-  typedef typename Tensor<float, 1>::DimensionPair DimensionPair;
+  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_evals<ColMajor>());
-  CALL_SUBTEST(test_scalar());
+  CALL_SUBTEST(test_evals<RowMajor>());
-  CALL_SUBTEST(test_multidims());
+  CALL_SUBTEST(test_scalar<ColMajor>());
-  CALL_SUBTEST(test_holes());
+  CALL_SUBTEST(test_scalar<RowMajor>());
-  CALL_SUBTEST(test_full_redux());
+  CALL_SUBTEST(test_multidims<ColMajor>());
-  CALL_SUBTEST(test_contraction_of_contraction());
+  CALL_SUBTEST(test_multidims<RowMajor>());
-  CALL_SUBTEST(test_expr());
+  CALL_SUBTEST(test_holes<ColMajor>());
-  CALL_SUBTEST(test_out_of_order_contraction());
+  CALL_SUBTEST(test_holes<RowMajor>());
-  CALL_SUBTEST(test_consistency());
+  CALL_SUBTEST(test_full_redux<ColMajor>());
-  CALL_SUBTEST(test_large_contraction());
+  CALL_SUBTEST(test_full_redux<RowMajor>());
-  CALL_SUBTEST(test_matrix_vector());
+  CALL_SUBTEST(test_contraction_of_contraction<ColMajor>());
-  CALL_SUBTEST(test_tensor_vector());
+  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>());
 }

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());
+}

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_(0,0) = 3.14f;
-    kernel_2d_(Eigen::array<int, 2>(1,0)) = 2.7f;
+    kernel_2d_(1,0) = 2.7f;
-    kernel_2d_(Eigen::array<int, 2>(0,1)) = 0.2f;
+    kernel_2d_(0,1) = 0.2f;
-    kernel_2d_(Eigen::array<int, 2>(1,1)) = 7.0f;
+    kernel_2d_(1,1) = 7.0f;
 
-    kernel_3d_(Eigen::array<int, 3>(0,0,0)) = 3.14f;
+    kernel_3d_(0,0,0) = 3.14f;
-    kernel_3d_(Eigen::array<int, 3>(0,1,0)) = 2.7f;
+    kernel_3d_(0,1,0) = 2.7f;
-    kernel_3d_(Eigen::array<int, 3>(0,0,1)) = 0.2f;
+    kernel_3d_(0,0,1) = 0.2f;
-    kernel_3d_(Eigen::array<int, 3>(0,1,1)) = 7.0f;
+    kernel_3d_(0,1,1) = 7.0f;
-    kernel_3d_(Eigen::array<int, 3>(1,0,0)) = -1.0f;
+    kernel_3d_(1,0,0) = -1.0f;
-    kernel_3d_(Eigen::array<int, 3>(1,1,0)) = -0.3f;
+    kernel_3d_(1,1,0) = -0.3f;
-    kernel_3d_(Eigen::array<int, 3>(1,0,1)) = -0.7f;
+    kernel_3d_(1,0,1) = -0.7f;
-    kernel_3d_(Eigen::array<int, 3>(1,1,1)) = -0.5f;
+    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, 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_, Eigen::array<int, 3>(2, 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> indices(0,0,0);
-  Eigen::DSizes<int, 3> sizes(Eigen::array<int, 3>(40,49,70));
+  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> indices(0,0,0);
-  Eigen::DSizes<int, 3> sizes(Eigen::array<int, 3>(40,49,69));
+  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> indices(0,0,0);
-  Eigen::DSizes<int, 3> sizes(Eigen::array<int, 3>(39,49,69));
+  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> in1(40,50,70);
-  Eigen::Tensor<float, 3> in2(Eigen::array<int, 3>(40,50,70));
+  Eigen::Tensor<float, 3> in2(40,50,70);
-  Eigen::Tensor<float, 3> out(Eigen::array<int, 3>(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 result = out(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) +
+        const float expected = (in1(i,j,k) * 3.14f + in1(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(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 result = out(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 +
+        const float expected = (in1(i,j,k) * 3.14f + in1(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(i,j,k+1) * 0.2f + in1(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(i+1,j,k) * -1.0f + in1(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);
+                                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> in1(40,50,70);
-  Eigen::Tensor<float, 3> in2(Eigen::array<int, 3>(40,50,70));
+  Eigen::Tensor<float, 3> in2(40,50,70);
-  Eigen::Tensor<float, 3> out(Eigen::array<int, 3>(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_in1(d_in1, 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_in2(d_in2, 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_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 result = out(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 +
+        const float expected = (in1(i,j,k) * 3.14f + in1(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(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 result = out(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 +
+        const float expected = (in1(i,j,k) * 3.14f + in1(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(i,j,k+1) * 0.2f + in1(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(i+1,j,k) * -1.0f + in1(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);
+                                in1(i+1,j,k+1) * -0.7f + in1(i+1,j+1,k+1) * -0.5f);
         VERIFY_IS_APPROX(expected, result);
       }
     }

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);
 }
 
 

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());

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());
 }

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) {
+              int row_offset = r*stride + i - row_padding;
-                expected = tensor(d, r-1+i, c-1+j, b);
+              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) {
+            int row_offset = r*stride + i - row_padding;
-              expected = tensor(d, r-1+i, c-1+j);
+            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());
 }

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);

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);
+  if (DataLayout == ColMajor) {
-  auto slice2 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice());
+    extents = Eigen::DSizes<ptrdiff_t, 4>(2,1,1,1);
-  VERIFY_IS_EQUAL(slice2.dimensions().TotalSize(), 2ul);
+    auto slice2 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice());
-  VERIFY_IS_EQUAL(slice2.data()[0], tensor(1,2,3,4));
+    VERIFY_IS_EQUAL(slice2.dimensions().TotalSize(), 2ul);
-  VERIFY_IS_EQUAL(slice2.data()[1], tensor(2,2,3,4));
+    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);
+  if (DataLayout == ColMajor) {
-  extents = Eigen::DSizes<ptrdiff_t, 4>(3,2,1,1);
+    offsets = Eigen::DSizes<ptrdiff_t, 4>(0,2,3,4);
-  auto slice4 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice());
+    extents = Eigen::DSizes<ptrdiff_t, 4>(3,2,1,1);
-  VERIFY_IS_EQUAL(slice4.dimensions().TotalSize(), 6ul);
+    auto slice4 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice());
-  for (int i = 0; i < 3; ++i) {
+    VERIFY_IS_EQUAL(slice4.dimensions().TotalSize(), 6ul);
-    for (int j = 0; j < 2; ++j) {
+    for (int i = 0; i < 3; ++i) {
-      VERIFY_IS_EQUAL(slice4.data()[i+3*j], tensor(i,2+j,3,4));
+      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);
+  if (DataLayout == ColMajor) {
-  extents = Eigen::DSizes<ptrdiff_t, 4>(3,5,7,2);
+    offsets = Eigen::DSizes<ptrdiff_t, 4>(0,0,0,4);
-  auto slice5 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice());
+    extents = Eigen::DSizes<ptrdiff_t, 4>(3,5,7,2);
-  VERIFY_IS_EQUAL(slice5.dimensions().TotalSize(), 210ul);
+    auto slice5 = SliceEvaluator(tensor.slice(offsets, extents), DefaultDevice());
-  for (int i = 0; i < 3; ++i) {
+    VERIFY_IS_EQUAL(slice5.dimensions().TotalSize(), 210ul);
-    for (int j = 0; j < 5; ++j) {
+    for (int i = 0; i < 3; ++i) {
-      for (int k = 0; k < 7; ++k) {
+      for (int j = 0; j < 5; ++j) {
-        for (int l = 0; l < 2; ++l) {
+        for (int k = 0; k < 7; ++k) {
-          int slice_index = i + 3 * (j + 5 * (k + 7 * l));
+          for (int l = 0; l < 2; ++l) {
-          VERIFY_IS_EQUAL(slice5.data()[slice_index], tensor(i,j,k,l+4));
+            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 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_as_lvalue<ColMajor>());
-  CALL_SUBTEST(test_slice_raw_data());
+  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());
 }

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];
+  std::string data1[6];
-  TensorMap<Tensor<string, 2>> mat1(data1, 2, 3);
+  TensorMap<Tensor<std::string, 2>> mat1(data1, 2, 3);
-  string data2[6];
+  std::string data2[6];
-  const TensorMap<Tensor<const string, 2>> mat2(data2, 2, 3);
+  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<std::string, 2> rslt3 = mat1;
-  Tensor<string, 2> rslt4 = mat2;
+  Tensor<std::string, 2> rslt4 = mat2;
 
-  Tensor<string, 2> rslt5(mat1);
+  Tensor<std::string, 2> rslt5(mat1);
-  Tensor<string, 2> rslt6(mat2);
+  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<std::string, 2> t1(2, 3);
-  Tensor<string, 2> t2(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<std::string, 2> t1 = data.slice(first_half, half_size);
-  Tensor<string, 2> t2 = data.slice(second_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<std::string, 1> data1(3);
-  Tensor<string, 1> data2(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());
 }

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_simple_padding<ColMajor>());
-  CALL_SUBTEST(test_padded_expr());
+  CALL_SUBTEST(test_simple_padding<RowMajor>());
+  CALL_SUBTEST(test_padded_expr<ColMajor>());
+  CALL_SUBTEST(test_padded_expr<RowMajor>());
 }

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;

unsupported/test/cxx11_tensor_reduction.cpp

@@ -13,15 +13,15 @@
 
 using Eigen::Tensor;
 
-static void test_simple_reductions()
+template <int DataLayout>
-{
+static void test_simple_reductions() {
-  Tensor<float, 4> tensor(2,3,5,7);
+  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);
+
+    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()
+static void test_full_reductions() {
-{
+  Tensor<float, 2, DataLayout> tensor(2, 3);
-  Tensor<float, 2> 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) {}
+  static const bool PacketAccess = false;
-  void reduce(const float val) {
+  UserReducer(float offset) : offset_(offset) {}
-    sum_ += val * val;
+  void reduce(const float val, float* accum) { *accum += val * val; }
-  }
+  float initialize() const { return 0; }
-  float finalize() const {
+  float finalize(const float accum) const { return 1.0f / (accum + offset_); }
-    return 1.0f / (sum_ + offset_);
-  }
 
  private:
-  float offset_;
+  const float offset_;
-  float sum_;
 };
 
-static void test_user_defined_reductions()
+template <int DataLayout>
-{
+static void test_user_defined_reductions() {
-  Tensor<float, 2> tensor(5,7);
+  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()
   }
 }
 
-
+template <int DataLayout>
-static void test_tensor_maps()
+static void test_tensor_maps() {
-{
+  int inputs[2 * 3 * 5 * 7];
-  int inputs[2*3*5*7];
+  TensorMap<Tensor<int, 4, DataLayout> > tensor_map(inputs, 2, 3, 5, 7);
-  TensorMap<Tensor<int, 4> > tensor_map(inputs, 2,3,5,7);
+  TensorMap<Tensor<const int, 4, DataLayout> > tensor_map_const(inputs, 2, 3, 5,
-  TensorMap<Tensor<const int, 4> > tensor_map_const(inputs, 2,3,5,7);
+                                                                7);
-  const TensorMap<Tensor<const int, 4> > tensor_map_const_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, DataLayout> result = tensor_map.sum(reduction_axis);
-  Tensor<int, 2> result2 = tensor_map_const.sum(reduction_axis);
+  Tensor<int, 2, DataLayout> result2 = tensor_map_const.sum(reduction_axis);
-  Tensor<int, 2> result3 = tensor_map_const_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();
 
-void test_cxx11_tensor_reduction()
+#if __cplusplus <= 199711L
-{
+  array<int, 2> reduction_axis;
-   CALL_SUBTEST(test_simple_reductions());
+  reduction_axis[0] = 1;
-   CALL_SUBTEST(test_full_reductions());
+  reduction_axis[1] = 3;
-   CALL_SUBTEST(test_user_defined_reductions());
+#else
-   CALL_SUBTEST(test_tensor_maps());
+  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<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>());
 }

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_simple_shuffling<ColMajor>());
-   CALL_SUBTEST(test_expr_shuffling());
+   CALL_SUBTEST(test_simple_shuffling<RowMajor>());
-   CALL_SUBTEST(test_shuffling_as_value());
+   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>());
 }

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);

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>());
 }

unsupported/test/cxx11_tensor_thread_pool.cpp

@@ -9,10 +9,10 @@
 
 #define EIGEN_USE_THREADS
 
-#include <iostream>
-#include "main.h"
-#include <Eigen/CXX11/Tensor>
 
+#include "main.h"
+#include <iostream>
+#include <Eigen/CXX11/Tensor>
 
 using Eigen::Tensor;
 
@@ -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, 4, DataLayout> t_left(30, 50, 37, 31);
-  Tensor<float, 5> t_right(37, 31, 70, 2, 10);
+  Tensor<float, 5, DataLayout> t_right(37, 31, 70, 2, 10);
-  Tensor<float, 5> t_result(30, 50, 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<Matrix<float, Dynamic, Dynamic, DataLayout>> MapXf;
-  typedef Map<MatrixXf> 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, DataLayout> t_left(32, 500);
-  Tensor<float, 2> t_right(32, 28*28);
+  Tensor<float, 2, DataLayout> t_right(32, 28*28);
-  Tensor<float, 2> t_result(500, 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),
+  Tensor<float, 3, DataLayout> left(internal::random<int>(1, 80),
-                        contract_size,
+                                    contract_size,
-                        internal::random<int>(1, 100));
+                                    internal::random<int>(1, 100));
 
-  Tensor<float, 4> right(internal::random<int>(1, 25),
+  Tensor<float, 4, DataLayout> right(internal::random<int>(1, 25),
-                         internal::random<int>(1, 37),
+                                     internal::random<int>(1, 37),
-                         contract_size,
+                                     contract_size,
-                         internal::random<int>(1, 51));
+                                     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());
 }