Merge pull request #67 from kokkos/batched-develop

Add cuda test and fix weird segmentation error coming from enum conversion.
kokkos · Sep 1, 2017 · c0932a7 · c0932a7
2 parents f2c730c + d7b9a92
commit c0932a7
Show file tree

Hide file tree

Showing 19 changed files with 61 additions and 25 deletions.
diff --git a/src/batched/KokkosBatched_Util.hpp b/src/batched/KokkosBatched_Util.hpp
@@ -30,18 +30,18 @@ namespace KokkosBatched {
 
     template<typename T>
     struct is_vector {
-      static constexpr bool value = false;
+      static const bool value = false;
     };
 
     template<typename Ta, typename Tb>
     struct is_same_mag_type {
-      static constexpr bool is_specialized = ( Kokkos::Details::ArithTraits<Ta>::is_specialized &&
+      static const bool is_specialized = ( Kokkos::Details::ArithTraits<Ta>::is_specialized &&
                                                Kokkos::Details::ArithTraits<Tb>::is_specialized );
 
-      static constexpr bool is_mag_type_same = std::is_same<typename Kokkos::Details::ArithTraits<Ta>::mag_type,
+      static const bool is_mag_type_same = std::is_same<typename Kokkos::Details::ArithTraits<Ta>::mag_type,
                                                             typename Kokkos::Details::ArithTraits<Tb>::mag_type>::value;
 
-      static constexpr bool value = is_specialized && is_mag_type_same;
+      static const bool value = is_specialized && is_mag_type_same;
     };
 
     // to use double, std::complex<double>, Kokkos::complex<double>
@@ -236,7 +236,7 @@ namespace KokkosBatched {
 		     "KokkosKernels:: Invalid VectorUnitTag<> type." );
 
       using type = VectorTag;
-      enum : int { length = l };
+      static const int length = l;
     };
 
 
@@ -258,8 +258,8 @@ namespace KokkosBatched {
     };
 
     struct Diag {
-      struct Unit    { enum : bool { use_unit_diag = true }; };
-      struct NonUnit { enum : bool { use_unit_diag = false}; };
+      struct Unit    { static const bool use_unit_diag = true;  };
+      struct NonUnit { static const bool use_unit_diag = false; };
     };
 
     struct Algo {

diff --git a/src/batched/KokkosBatched_Vector.hpp b/src/batched/KokkosBatched_Vector.hpp
@@ -12,12 +12,12 @@ namespace KokkosBatched {
 
     template<typename T, typename SpT, int l>
     struct is_vector<Vector<VectorTag<SIMD<T,SpT>,l> > > {
-      static constexpr bool value = true;
+      static const bool value = true;
     };
 
     template<typename T, typename SpT, int l>
     struct is_vector<Vector<VectorTag<AVX<T,SpT>,l> > > {
-      static constexpr bool value = true;
+      static const bool value = true;
     };
 
   }

diff --git a/src/batched/KokkosBatched_Vector_AVX256D.hpp b/src/batched/KokkosBatched_Vector_AVX256D.hpp
@@ -21,7 +21,7 @@ namespace KokkosBatched {
       using value_type = double;
       using mag_type = double;
 
-      enum : int { vector_length = 4 };
+      static const int vector_length = 4;
 
       union data_type {
         __m256d v;
@@ -278,7 +278,7 @@ namespace KokkosBatched {
       using value_type = Kokkos::complex<double>;
       using mag_type = double;
 
-      enum : int { vector_length = 2 };
+      static const int vector_length = 2;
 
       union data_type {
         __m256d v;

diff --git a/src/batched/KokkosBatched_Vector_AVX512D.hpp b/src/batched/KokkosBatched_Vector_AVX512D.hpp
@@ -19,7 +19,7 @@ namespace KokkosBatched {
       using value_type = double;
       using mag_type = double;
 
-      enum : int { vector_length = 8 };
+      static const int vector_length = 8;
 
       union data_type {
         __m512d v;
@@ -278,7 +278,7 @@ namespace KokkosBatched {
       using value_type = Kokkos::complex<double>;
       using mag_type = double;
 
-      enum : int { vector_length = 4 };
+      static const int vector_length = 4;
 
       union data_type {
         __m512d v;

diff --git a/src/batched/KokkosBatched_Vector_SIMD.hpp b/src/batched/KokkosBatched_Vector_SIMD.hpp
@@ -22,7 +22,7 @@ namespace KokkosBatched {
       //   - for host, the lambda does not allow further compiler optimization 
       using member_type = typename tag_type::member_type;
 
-      enum : int { vector_length = tag_type::length };
+      static const int vector_length = tag_type::length;
 
       typedef value_type data_type[vector_length];
 
@@ -87,7 +87,7 @@ namespace KokkosBatched {
         Kokkos::parallel_for
           (Kokkos::Impl::ThreadVectorRangeBoundariesStruct<int,member_type>(vector_length),
            [&](const int &i) {
-
+            _data[i] = b._data[i];
           });
 #else
 #if defined( KOKKOS_ENABLE_PRAGMA_IVDEP )

diff --git a/unit_test/Makefile b/unit_test/Makefile
@@ -181,6 +181,18 @@ ifeq ($(KOKKOSKERNELS_INTERNAL_TEST_CUDA), 1)
   OBJ_CUDA += Test_Cuda_Graph_graph_color.o
   OBJ_CUDA += Test_Cuda_Common_ArithTraits.o
   OBJ_CUDA += Test_Cuda_Common_set_bit_count.o
+  OBJ_CUDA += Test_Cuda_Batched_SerialMatUtil.o
+  OBJ_CUDA += Test_Cuda_Batched_SerialGemm.o
+  OBJ_CUDA += Test_Cuda_Batched_SerialTrsm.o
+  OBJ_CUDA += Test_Cuda_Batched_SerialLU.o
+  OBJ_CUDA += Test_Cuda_Batched_SerialGemv.o
+  OBJ_CUDA += Test_Cuda_Batched_SerialTrsv.o
+  OBJ_CUDA += Test_Cuda_Batched_TeamMatUtil.o
+  OBJ_CUDA += Test_Cuda_Batched_TeamGemm.o
+  OBJ_CUDA += Test_Cuda_Batched_TeamTrsm.o
+  OBJ_CUDA += Test_Cuda_Batched_TeamLU.o
+  OBJ_CUDA += Test_Cuda_Batched_TeamGemv.o
+  OBJ_CUDA += Test_Cuda_Batched_TeamTrsv.o
   TARGETS += KokkosKernels_UnitTest_Cuda
   TEST_TARGETS += test-cuda
 endif

diff --git a/unit_test/batched/Test_Batched_VectorArithmatic.hpp b/unit_test/batched/Test_Batched_VectorArithmatic.hpp
@@ -17,7 +17,7 @@ namespace Test {
     /// random data initialization
     typedef Vector<VectorTagType> vector_type;
     typedef typename vector_type::value_type value_type;    
-    constexpr int vector_length = vector_type::vector_length;
+    const int vector_length = vector_type::vector_length;
 
     typedef Kokkos::Details::ArithTraits<value_type> ats;
     typedef typename ats::mag_type mag_type;
@@ -171,21 +171,21 @@ int test_batched_vector_arithmatic() {
 
 #if defined(KOKKOSKERNELS_INST_FLOAT)
 TEST_F( TestCategory, batched_vector_arithmatic_simd_float8 ) {
-  typedef VectorTag<SIMD<float>, 8> vector_tag_type;
+  typedef VectorTag<SIMD<float,TestExecSpace>, 8> vector_tag_type;
   test_batched_vector_arithmatic<TestExecSpace,vector_tag_type>();
 }
 #endif
 
 #if defined(KOKKOSKERNELS_INST_DOUBLE)
 TEST_F( TestCategory, batched_vector_arithmatic_simd_double4 ) {
-  typedef VectorTag<SIMD<double>, 4> vector_tag_type;
+  typedef VectorTag<SIMD<double,TestExecSpace>, 4> vector_tag_type;
   test_batched_vector_arithmatic<TestExecSpace,vector_tag_type>();
 }
 #endif
 
 #if defined(KOKKOSKERNELS_INST_COMPLEX_DOUBLE)
 TEST_F( TestCategory, batched_vector_arithmatic_simd_dcomplex2 ) {
-  typedef VectorTag<SIMD<Kokkos::complex<double> >, 2> vector_tag_type;
+  typedef VectorTag<SIMD<Kokkos::complex<double>,TestExecSpace>, 2> vector_tag_type;
   test_batched_vector_arithmatic<TestExecSpace,vector_tag_type>();
 }
 #endif
@@ -197,21 +197,21 @@ TEST_F( TestCategory, batched_vector_arithmatic_simd_dcomplex2 ) {
 #if defined(__AVX__) || defined(__AVX2__)
 #if defined(KOKKOSKERNELS_INST_FLOAT)
 TEST_F( TestCategory, batched_vector_arithmatic_avx_float8 ) {
-  typedef VectorTag<AVX<float>, 8> vector_tag_type;
+  typedef VectorTag<AVX<float,TestExecSpace>, 8> vector_tag_type;
   test_batched_vector_arithmatic<TestExecSpace,vector_tag_type>();
 }
 #endif
 
 #if defined(KOKKOSKERNELS_INST_DOUBLE)
 TEST_F( TestCategory, batched_vector_arithmatic_avx_double4 ) {
-  typedef VectorTag<AVX<double>, 4> vector_tag_type;
+  typedef VectorTag<AVX<double,TestExecSpace>, 4> vector_tag_type;
   test_batched_vector_arithmatic<TestExecSpace,vector_tag_type>();
 }
 #endif
 
 #if defined(KOKKOSKERNELS_INST_COMPLEX_DOUBLE)
 TEST_F( TestCategory, batched_vector_arithmatic_avx_dcomplex2 ) {
-  typedef VectorTag<AVX<Kokkos::complex<double> >, 2> vector_tag_type;
+  typedef VectorTag<AVX<Kokkos::complex<double>,TestExecSpace>, 2> vector_tag_type;
   test_batched_vector_arithmatic<TestExecSpace,vector_tag_type>();
 }
 #endif
@@ -224,21 +224,21 @@ TEST_F( TestCategory, batched_vector_arithmatic_avx_dcomplex2 ) {
 #if defined(__AVX512F__)
 #if defined(KOKKOSKERNELS_INST_FLOAT)
 TEST_F( TestCategory, batched_vector_arithmatic_avx_float16 ) {
-  typedef VectorTag<AVX<float>, 16> vector_tag_type;
+  typedef VectorTag<AVX<float,TestExecSpace>, 16> vector_tag_type;
   test_batched_vector_arithmatic<TestExecSpace,vector_tag_type>();
 }
 #endif
 
 #if defined(KOKKOSKERNELS_INST_DOUBLE)
 TEST_F( TestCategory, batched_vector_arithmatic_avx_double8 ) {
-  typedef VectorTag<AVX<double>, 8> vector_tag_type;
+  typedef VectorTag<AVX<double,TestExecSpace>, 8> vector_tag_type;
   test_batched_vector_arithmatic<TestExecSpace,vector_tag_type>();
 }
 #endif
 
 #if defined(KOKKOSKERNELS_INST_COMPLEX_DOUBLE)
 TEST_F( TestCategory, batched_vector_arithmatic_avx_dcomplex4 ) {
-  typedef VectorTag<AVX<Kokkos::complex<double> >, 4> vector_tag_type;
+  typedef VectorTag<AVX<Kokkos::complex<double>,TestExecSpace>, 4> vector_tag_type;
   test_batched_vector_arithmatic<TestExecSpace,vector_tag_type>();
 }
 #endif

diff --git a/unit_test/cuda/Test_Cuda_Batched_SerialGemm.cpp b/unit_test/cuda/Test_Cuda_Batched_SerialGemm.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_SerialGemm.hpp"
diff --git a/unit_test/cuda/Test_Cuda_Batched_SerialGemv.cpp b/unit_test/cuda/Test_Cuda_Batched_SerialGemv.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_SerialGemv.hpp"
diff --git a/unit_test/cuda/Test_Cuda_Batched_SerialLU.cpp b/unit_test/cuda/Test_Cuda_Batched_SerialLU.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_SerialLU.hpp"
diff --git a/unit_test/cuda/Test_Cuda_Batched_SerialMatUtil.cpp b/unit_test/cuda/Test_Cuda_Batched_SerialMatUtil.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_SerialMatUtil.hpp"
diff --git a/unit_test/cuda/Test_Cuda_Batched_SerialTrsm.cpp b/unit_test/cuda/Test_Cuda_Batched_SerialTrsm.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_SerialTrsm.hpp"
diff --git a/unit_test/cuda/Test_Cuda_Batched_SerialTrsv.cpp b/unit_test/cuda/Test_Cuda_Batched_SerialTrsv.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_SerialTrsv.hpp"
diff --git a/unit_test/cuda/Test_Cuda_Batched_TeamGemm.cpp b/unit_test/cuda/Test_Cuda_Batched_TeamGemm.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_TeamGemm.hpp"
diff --git a/unit_test/cuda/Test_Cuda_Batched_TeamGemv.cpp b/unit_test/cuda/Test_Cuda_Batched_TeamGemv.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_TeamGemv.hpp"
diff --git a/unit_test/cuda/Test_Cuda_Batched_TeamLU.cpp b/unit_test/cuda/Test_Cuda_Batched_TeamLU.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_TeamLU.hpp"
diff --git a/unit_test/cuda/Test_Cuda_Batched_TeamMatUtil.cpp b/unit_test/cuda/Test_Cuda_Batched_TeamMatUtil.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_TeamMatUtil.hpp"
diff --git a/unit_test/cuda/Test_Cuda_Batched_TeamTrsm.cpp b/unit_test/cuda/Test_Cuda_Batched_TeamTrsm.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_TeamTrsm.hpp"
diff --git a/unit_test/cuda/Test_Cuda_Batched_TeamTrsv.cpp b/unit_test/cuda/Test_Cuda_Batched_TeamTrsv.cpp
@@ -0,0 +1,2 @@
+#include "Test_Cuda.hpp"
+#include "Test_Batched_TeamTrsv.hpp"
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		#include "Test_Cuda.hpp"
		#include "Test_Batched_SerialGemm.hpp"
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		#include "Test_Cuda.hpp"
		#include "Test_Batched_SerialGemv.hpp"
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		#include "Test_Cuda.hpp"
		#include "Test_Batched_SerialLU.hpp"
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		#include "Test_Cuda.hpp"
		#include "Test_Batched_SerialMatUtil.hpp"
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		#include "Test_Cuda.hpp"
		#include "Test_Batched_SerialTrsm.hpp"
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		#include "Test_Cuda.hpp"
		#include "Test_Batched_SerialTrsv.hpp"
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		#include "Test_Cuda.hpp"
		#include "Test_Batched_TeamGemm.hpp"
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		#include "Test_Cuda.hpp"
		#include "Test_Batched_TeamGemv.hpp"
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		#include "Test_Cuda.hpp"
		#include "Test_Batched_TeamLU.hpp"
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		#include "Test_Cuda.hpp"
		#include "Test_Batched_TeamMatUtil.hpp"