Merge pull request #719 from kokkos/release-3.1.1

Release 3.1.1

CHANGELOG.md

@@ -1,6 +1,13 @@
 # Change Log
 
 
+## [3.1.1](https://github.com/kokkos/kokkos-kernels/tree/3.1.1) (2020-05-04)
+[Full Changelog](https://github.com/kokkos/kokkos-kernels/compare/3.1.00...3.1.1)
+
+** Fixed bugs:** 
+
+- KokkosBatched QR PR breaking nightly tests [\#691](https://github.com/kokkos/kokkos-kernels/issues/691)
+
 ## [3.1.00](https://github.com/kokkos/kokkos-kernels/tree/3.1.00) (2020-04-14)
 [Full Changelog](https://github.com/kokkos/kokkos-kernels/compare/3.0.00...3.1.00)
 

CMakeLists.txt

@@ -28,7 +28,7 @@ IF(NOT KOKKOSKERNELS_HAS_TRILINOS)
   ENDIF()
   SET(KokkosKernels_VERSION_MAJOR 3)
   SET(KokkosKernels_VERSION_MINOR 1)
-  SET(KokkosKernels_VERSION_PATCH 0)
+  SET(KokkosKernels_VERSION_PATCH 1)
 ENDIF()
 
 IF(${CMAKE_VERSION} VERSION_GREATER_EQUAL "3.12.0") 

master_history.txt

@@ -7,4 +7,5 @@ tag: 2.7.24     date: 11/05/2018  master: 1a7b524b    develop: fab89e37
 tag: 2.8.00     date: 02/05/2019  master: a6e05e06    develop: 6a790321
 tag: 2.9.00     date: 06/24/2019  master: 4ee5f3c6    develop: 094da30c
 tag: 3.0.00     date: 01/31/2020  master: d86db111    release-candidate-3.0: cf24ab90
-tag: 3.1.00     date: 04/14/2020  master: f199f45d    develop: 8d063eae 
+tag: 3.1.00     date: 04/14/2020  master: f199f45d    develop: 8d063eae
+tag: 3.1.1      date: 05/04/2020  master: 43773523    release: 6fce7502 

perf_test/sparse/KokkosSparse_sptrsv_superlu.cpp

@@ -702,7 +702,7 @@ int main(int argc, char **argv) {
       using scalar_t = double;
       scalarTypeString = "(scalar_t = double)";
     #else
-      #error "Invalid type specified in KOKKOSKERNELS_SCALARS, supported types are "double,complex<double>""
+      #error "Invalid type specified in KOKKOSKERNELS_SCALARS, supported types are \"double,complex<double>\""
     #endif
   #endif
   int total_errors = test_sptrsv_perf<scalar_t> (tests, verbose, filename, symm_mode, metis, merge,

src/batched/KokkosBatched_Vector_SIMD.hpp

@@ -45,7 +45,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)
@@ -59,7 +59,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)
@@ -73,7 +73,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)
@@ -88,7 +88,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)
@@ -109,7 +109,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)
@@ -529,7 +529,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)
@@ -546,7 +546,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)
@@ -639,7 +639,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)
@@ -723,7 +723,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)
@@ -739,7 +739,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)
@@ -824,7 +824,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)
@@ -840,7 +840,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__)
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
       for (int i=0;i<vector_length;++i)

src/batched/KokkosBatched_Vector_SIMD_Arith.hpp

@@ -4,6 +4,7 @@
 /// \author Kyungjoo Kim ([email protected])
 
 #include "Kokkos_Complex.hpp"
+#include "KokkosKernels_Macros.hpp"
 
 namespace KokkosBatched {
 
@@ -63,7 +64,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__) 
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
     for (int i=0;i<l;++i)
@@ -279,7 +280,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__) 
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
     for (int i=0;i<l;++i)
@@ -342,7 +343,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__) 
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
     for (int i=0;i<l;++i)
@@ -539,7 +540,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__) 
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
     for (int i=0;i<l;++i)
@@ -824,7 +825,7 @@ namespace KokkosBatched {
 #if defined( KOKKOS_ENABLE_PRAGMA_VECTOR )
 #pragma vector always
 #endif
-#if defined( KOKKOS_ENABLE_OPENMP ) && !defined(__CUDA_ARCH__) 
+#ifdef KOKKOSKERNELS_ENABLE_OMP_SIMD
 #pragma omp simd
 #endif
     for (int i=0;i<l;++i)

src/blas/impl/KokkosBlas3_gemm_impl.hpp

@@ -45,7 +45,8 @@
 #ifndef KOKKOS_BLAS3_GEMM_IMPL_HPP_
 #define KOKKOS_BLAS3_GEMM_IMPL_HPP_
 
-#include<Kokkos_Core.hpp>
+#include <Kokkos_Core.hpp>
+#include "KokkosKernels_Macros.hpp"
 
 #ifdef KOKKOS_ENABLE_CXX14
 #ifdef KOKKOS_COMPILER_GNU
@@ -401,16 +402,10 @@ void impl_team_gemm_block(const TeamHandle& team, const ViewTypeC& C, const View
   const int blockB1 = B.extent_int(1);
 #endif
   Kokkos::parallel_for(Kokkos::TeamThreadRange(team,blockA0), [&] (const int i) {
-#if defined(__CUDA_ARCH__) || !defined(KOKKOS_ENABLE_OPENMP)
+#ifndef KOKKOSKERNELS_ENABLE_OMP_SIMD
     Kokkos::parallel_for(Kokkos::ThreadVectorRange(team,blockB1/4), [&] (const int B_j) {
 #else
-  #if defined(KOKKOS_COMPILER_GNU)
-    #if (KOKKOS_COMPILER_GNU > 485 )
     #pragma omp simd
-    #endif
-  #else
-    #pragma omp simd
-  #endif
     for(int B_j=0; B_j<blockB1/4; B_j++) {
 #endif
       ScalarC C_ij0 = 0;
@@ -428,7 +423,7 @@ void impl_team_gemm_block(const TeamHandle& team, const ViewTypeC& C, const View
       C(i,B_j+blockB1/4) += C_ij1;
       C(i,B_j+2*blockB1/4) += C_ij2;
       C(i,B_j+3*blockB1/4) += C_ij3;
-#if defined(__CUDA_ARCH__) || !defined(KOKKOS_ENABLE_OPENMP)
+#ifndef KOKKOSKERNELS_ENABLE_OMP_SIMD
     });
 #else
     }

src/common/KokkosKernels_Macros.hpp

@@ -0,0 +1,65 @@
+/*
+//@HEADER
+// ************************************************************************
+//
+//                        Kokkos v. 3.0
+//       Copyright (2020) National Technology & Engineering
+//               Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are
+// met:
+//
+// 1. Redistributions of source code must retain the above copyright
+// notice, this list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright
+// notice, this list of conditions and the following disclaimer in the
+// documentation and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the Corporation nor the names of the
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY NTESS "AS IS" AND ANY
+// EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+// PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NTESS OR THE
+// CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+// EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+// PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
+// LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
+// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+//
+// Questions? Contact Siva Rajamanickam ([email protected])
+//
+// ************************************************************************
+//@HEADER
+*/
+
+#ifndef _KOKKOSKERNELS_MACROUTILS_HPP_
+#define _KOKKOSKERNELS_MACROUTILS_HPP_
+
+// If KOKKOSKERNELS_ENABLE_OMP_SIMD is defined, it's legal to place
+// "#pragma omp simd" before a for loop. It's never defined if CUDA is enabled,
+// since in that case, Kokkos::ThreadVectorRange should be used instead for SIMD parallel loops.
+
+#if !defined(KOKKOS_ENABLE_CUDA) && defined(KOKKOS_ENABLE_OPENMP)
+  #if defined(KOKKOS_COMPILER_GNU)
+    // GCC 4.8.5 and older do not support #pragma omp simd
+    #if (KOKKOS_COMPILER_GNU > 485 )
+      #define KOKKOSKERNELS_ENABLE_OMP_SIMD
+    #endif
+  #else
+    // All other Kokkos-supported compilers support it.
+    #define KOKKOSKERNELS_ENABLE_OMP_SIMD
+  #endif
+#endif
+
+#endif //_KOKKOSKERNELS_MACROUTILS_HPP_
+

test_common/Test_Common_Sorting.hpp

@@ -569,24 +569,34 @@ void testSortCRS(default_lno_t numRows, default_size_type nnz, bool doValues)
   Kokkos::deep_copy(rowmapHost, rowmap);
   Kokkos::deep_copy(entriesHost, entries);
   Kokkos::deep_copy(valuesHost, values);
+  struct ColValue
+  {
+    ColValue() {}
+    ColValue(lno_t c, scalar_t v) : col(c), val(v) {}
+    bool operator<(const ColValue& rhs) const
+    {
+      return col < rhs.col;
+    }
+    bool operator==(const ColValue& rhs) const
+    {
+      return col == rhs.col && val == rhs.val;
+    }
+    lno_t col;
+    scalar_t val;
+  };
   //sort one row at a time on host using STL.
   {
-    using ColValue = std::pair<lno_t, scalar_t>;
     for(lno_t i = 0; i < numRows; i++)
     {
       std::vector<ColValue> rowCopy;
       for(size_type j = rowmapHost(i); j < rowmapHost(i + 1); j++)
         rowCopy.emplace_back(entriesHost(j), valuesHost(j));
-      std::sort(rowCopy.begin(), rowCopy.end(),
-          [](const ColValue& lhs, const ColValue& rhs)
-          {
-            return lhs.first < rhs.first;
-          });
+      std::sort(rowCopy.begin(), rowCopy.end());
       //write sorted row back
       for(size_t j = 0; j < rowCopy.size(); j++)
       {
-        entriesHost(rowmapHost(i) + j) = rowCopy[j].first;
-        valuesHost(rowmapHost(i) + j) = rowCopy[j].second;
+        entriesHost(rowmapHost(i) + j) = rowCopy[j].col;
+        valuesHost(rowmapHost(i) + j) = rowCopy[j].val;
       }
     }
   }