kokkos · crtrott · Feb 12, 2024 · Nov 22, 2023 · Nov 22, 2023 · Nov 22, 2023
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -1,5 +1,14 @@
 # Change Log
 
+## [4.2.01](https://github.com/kokkos/kokkos-kernels/tree/4.2.01) (2024-01-17)
+[Full Changelog](https://github.com/kokkos/kokkos-kernels/compare/4.2.00...4.2.01)
+
+### Bug Fixes:
+
+- LAPACK: magma tpl fixes [\#2044](https://github.com/kokkos/kokkos-kernels/pull/2044)
+- BLAS: fix bug in TPL layer of `KokkosBlas::swap` [\#2052](https://github.com/kokkos/kokkos-kernels/pull/2052)
+- ROCm 6 deprecation fixes for rocsparse [\#2050](https://github.com/kokkos/kokkos-kernels/pull/2050)
+
 ## [4.2.00](https://github.com/kokkos/kokkos-kernels/tree/4.2.00) (2023-11-06)
 [Full Changelog](https://github.com/kokkos/kokkos-kernels/compare/4.1.00...4.2.00)
 

diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -11,7 +11,7 @@ SET(KOKKOSKERNELS_TOP_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR})
 
 SET(KokkosKernels_VERSION_MAJOR 4)
 SET(KokkosKernels_VERSION_MINOR 2)
-SET(KokkosKernels_VERSION_PATCH 00)
+SET(KokkosKernels_VERSION_PATCH 1)
 SET(KokkosKernels_VERSION "${KokkosKernels_VERSION_MAJOR}.${KokkosKernels_VERSION_MINOR}.${KokkosKernels_VERSION_PATCH}")
 
 #Set variables for config file
@@ -133,7 +133,7 @@ ELSE()
         MESSAGE(WARNING "Configuring with Kokkos ${Kokkos_VERSION} which is newer than the expected develop branch - version check may need update")
       ENDIF()
     ELSE()
-      MESSAGE(FATAL_ERROR "Kokkos Kernels ${KokkosKernels_VERSION} requires 4.1.00, 4.2.00 or develop")
+      MESSAGE(FATAL_ERROR "Kokkos Kernels ${KokkosKernels_VERSION} requires 4.1.00, 4.2.00, 4.2.01 or develop")
     ENDIF()
   ENDIF()
 

diff --git a/blas/tpls/KokkosBlas1_swap_tpl_spec_decl.hpp b/blas/tpls/KokkosBlas1_swap_tpl_spec_decl.hpp
@@ -293,10 +293,10 @@ namespace Impl {
                            Kokkos::Device<EXECSPACE, MEMSPACE>,             \
                            Kokkos::MemoryTraits<Kokkos::Unmanaged>>,        \
               true, ETI_SPEC_AVAIL> {                                       \
-    using XVector = Kokkos::View<Kokkos::complex<float>, LAYOUT,            \
+    using XVector = Kokkos::View<Kokkos::complex<float>*, LAYOUT,           \
                                  Kokkos::Device<EXECSPACE, MEMSPACE>,       \
                                  Kokkos::MemoryTraits<Kokkos::Unmanaged>>;  \
-    using YVector = Kokkos::View<Kokkos::complex<float>, LAYOUT,            \
+    using YVector = Kokkos::View<Kokkos::complex<float>*, LAYOUT,           \
                                  Kokkos::Device<EXECSPACE, MEMSPACE>,       \
                                  Kokkos::MemoryTraits<Kokkos::Unmanaged>>;  \
     static void swap(EXECSPACE const& space, XVector const& X,              \

diff --git a/lapack/CMakeLists.txt b/lapack/CMakeLists.txt
@@ -28,7 +28,7 @@ IF (KOKKOSKERNELS_ENABLE_TPL_LAPACK OR KOKKOSKERNELS_ENABLE_TPL_MKL OR KOKKOSKER
 ENDIF()
 
 # Include cuda lapack TPL source file
-IF (KOKKOSKERNELS_ENABLE_TPL_CUSOLVER)
+IF (KOKKOSKERNELS_ENABLE_TPL_MAGMA)
   LIST(APPEND SOURCES
    lapack/tpls/KokkosLapack_Cuda_tpl.cpp
   )

diff --git a/lapack/tpls/KokkosLapack_Cuda_tpl.hpp b/lapack/tpls/KokkosLapack_Cuda_tpl.hpp
@@ -16,31 +16,8 @@
 #ifndef KOKKOSLAPACK_CUDA_TPL_HPP_
 #define KOKKOSLAPACK_CUDA_TPL_HPP_
 
-#if defined(KOKKOSKERNELS_ENABLE_TPL_CUSOLVER)
-#include <KokkosLapack_tpl_spec.hpp>
-
-namespace KokkosLapack {
-namespace Impl {
-
-CudaLapackSingleton::CudaLapackSingleton() {
-  cusolverStatus_t stat = cusolverDnCreate(&handle);
-  if (stat != CUSOLVER_STATUS_SUCCESS)
-    Kokkos::abort("CUSOLVER initialization failed\n");
-
-  Kokkos::push_finalize_hook([&]() { cusolverDnDestroy(handle); });
-}
-
-CudaLapackSingleton& CudaLapackSingleton::singleton() {
-  static CudaLapackSingleton s;
-  return s;
-}
-
-}  // namespace Impl
-}  // namespace KokkosLapack
-#endif  // defined (KOKKOSKERNELS_ENABLE_TPL_CUSOLVER)
-
 #if defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA)
-#include <KokkosLapack_tpl_spec.hpp>
+#include <KokkosLapack_magma.hpp>
 
 namespace KokkosLapack {
 namespace Impl {

diff --git a/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp b/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp
@@ -275,7 +275,7 @@ KOKKOSLAPACK_CGESV_LAPACK(Kokkos::LayoutLeft, Kokkos::HostSpace, false)
 
 // MAGMA
 #ifdef KOKKOSKERNELS_ENABLE_TPL_MAGMA
-#include <KokkosLapack_tpl_spec.hpp>
+#include <KokkosLapack_magma.hpp>
 
 namespace KokkosLapack {
 namespace Impl {

diff --git a/lapack/tpls/KokkosLapack_magma.hpp b/lapack/tpls/KokkosLapack_magma.hpp
@@ -0,0 +1,35 @@
+//@HEADER
+// ************************************************************************
+//
+//                        Kokkos v. 4.0
+//       Copyright (2022) 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.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//@HEADER
+
+#ifndef KOKKOSLAPACK_MAGMA_HPP_
+#define KOKKOSLAPACK_MAGMA_HPP_
+// If LAPACK TPL is enabled, it is preferred over magma's LAPACK
+#ifdef KOKKOSKERNELS_ENABLE_TPL_MAGMA
+#include "magma_v2.h"
+
+namespace KokkosLapack {
+namespace Impl {
+
+struct MagmaSingleton {
+  MagmaSingleton();
+
+  static MagmaSingleton& singleton();
+};
+
+}  // namespace Impl
+}  // namespace KokkosLapack
+#endif  // KOKKOSKERNELS_ENABLE_TPL_MAGMA
+#endif  // KOKKOSLAPACK_MAGMA_HPP_
diff --git a/lapack/tpls/KokkosLapack_trtri_tpl_spec_decl.hpp b/lapack/tpls/KokkosLapack_trtri_tpl_spec_decl.hpp
@@ -18,7 +18,9 @@
 #define KOKKOSLAPACK_TRTRI_TPL_SPEC_DECL_HPP_
 
 #include "KokkosLapack_Host_tpl.hpp"  // trtri prototype
-//#include "KokkosLapack_tpl_spec.hpp"
+#ifdef KOKKOSKERNELS_ENABLE_TPL_MAGMA
+#include "KokkosLapack_magma.hpp"
+#endif
 
 namespace KokkosLapack {
 namespace Impl {

diff --git a/master_history.txt b/master_history.txt
@@ -23,3 +23,4 @@ tag: 4.0.00     date: 02/23/2023  master: b4014bf2    release: a10dff20
 tag: 4.0.01     date: 04/26/2023  master: b9c1bab7    release: 8809e41c
 tag: 4.1.00     date: 06/20/2023  master: 1331baf1    release: 14ad220a
 tag: 4.2.00     date: 11/09/2023  master: 25a31f88    release: 912d3778
+tag: 4.2.01     date: 01/30/2024  master: f429f6ec    release: bcf9854b
diff --git a/sparse/src/KokkosSparse_Utils_rocsparse.hpp b/sparse/src/KokkosSparse_Utils_rocsparse.hpp
@@ -21,8 +21,12 @@
 #include <sstream>
 
 #ifdef KOKKOSKERNELS_ENABLE_TPL_ROCSPARSE
+#if __has_include(<rocm-core/rocm_version.h>)
+#include <rocm-core/rocm_version.h>
+#else
 #include <rocm_version.h>
-#include "rocsparse/rocsparse.h"
+#endif
+#include <rocsparse/rocsparse.h>
 
 namespace KokkosSparse {
 namespace Impl {

diff --git a/sparse/tpls/KokkosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp b/sparse/tpls/KokkosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp
@@ -869,8 +869,46 @@ void spmv_block_impl_rocsparse(
   rocsparse_mat_info info;
   KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_create_mat_info(&info));
 
+  // *_ex* functions deprecated in introduced in 6+
+#if KOKKOSSPARSE_IMPL_ROCM_VERSION >= 60000
+  if constexpr (std::is_same_v<value_type, float>) {
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_sbsrmv_analysis(
+        handle, dir, trans, mb, nb, nnzb, descr, bsr_val, bsr_row_ptr,
+        bsr_col_ind, block_dim, info));
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_sbsrmv(
+        handle, dir, trans, mb, nb, nnzb, alpha_, descr, bsr_val, bsr_row_ptr,
+        bsr_col_ind, block_dim, info, x_, beta_, y_));
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_bsrsv_clear(handle, info));
+  } else if constexpr (std::is_same_v<value_type, double>) {
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_dbsrmv_analysis(
+        handle, dir, trans, mb, nb, nnzb, descr, bsr_val, bsr_row_ptr,
+        bsr_col_ind, block_dim, info));
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_dbsrmv(
+        handle, dir, trans, mb, nb, nnzb, alpha_, descr, bsr_val, bsr_row_ptr,
+        bsr_col_ind, block_dim, info, x_, beta_, y_));
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_bsrsv_clear(handle, info));
+  } else if constexpr (std::is_same_v<value_type, Kokkos::complex<float>>) {
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_cbsrmv_analysis(
+        handle, dir, trans, mb, nb, nnzb, descr, bsr_val, bsr_row_ptr,
+        bsr_col_ind, block_dim, info));
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_cbsrmv(
+        handle, dir, trans, mb, nb, nnzb, alpha_, descr, bsr_val, bsr_row_ptr,
+        bsr_col_ind, block_dim, info, x_, beta_, y_));
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_bsrsv_clear(handle, info));
+  } else if constexpr (std::is_same_v<value_type, Kokkos::complex<double>>) {
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_zbsrmv_analysis(
+        handle, dir, trans, mb, nb, nnzb, descr, bsr_val, bsr_row_ptr,
+        bsr_col_ind, block_dim, info));
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_zbsrmv(
+        handle, dir, trans, mb, nb, nnzb, alpha_, descr, bsr_val, bsr_row_ptr,
+        bsr_col_ind, block_dim, info, x_, beta_, y_));
+    KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_bsrsv_clear(handle, info));
+  } else {
+    static_assert(KokkosKernels::Impl::always_false_v<value_type>,
+                  "unsupported value type for rocsparse_*bsrmv");
+  }
   // *_ex* functions introduced in 5.4.0
-#if KOKKOSSPARSE_IMPL_ROCM_VERSION < 50400
+#elif KOKKOSSPARSE_IMPL_ROCM_VERSION < 50400
   if constexpr (std::is_same_v<value_type, float>) {
     KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_sbsrmv(
         handle, dir, trans, mb, nb, nnzb, alpha_, descr, bsr_val, bsr_row_ptr,

diff --git a/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp b/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp
@@ -357,8 +357,6 @@ KOKKOSSPARSE_SPMV_CUSPARSE(Kokkos::complex<float>, int64_t, size_t,
 
 // rocSPARSE
 #if defined(KOKKOSKERNELS_ENABLE_TPL_ROCSPARSE)
-#include <rocsparse/rocsparse.h>
-#include <rocm_version.h>
 #include "KokkosSparse_Utils_rocsparse.hpp"
 
 namespace KokkosSparse {
@@ -441,7 +439,17 @@ void spmv_rocsparse(const Kokkos::HIP& exec,
       alg = rocsparse_spmv_alg_csr_stream;
   }
 
-#if KOKKOSSPARSE_IMPL_ROCM_VERSION >= 50400
+#if KOKKOSSPARSE_IMPL_ROCM_VERSION >= 60000
+  KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(
+      rocsparse_spmv(handle, myRocsparseOperation, &alpha, Aspmat, vecX, &beta,
+                     vecY, compute_type, alg, rocsparse_spmv_stage_buffer_size,
+                     &buffer_size, tmp_buffer));
+  KOKKOS_IMPL_HIP_SAFE_CALL(hipMalloc(&tmp_buffer, buffer_size));
+  KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(
+      rocsparse_spmv(handle, myRocsparseOperation, &alpha, Aspmat, vecX, &beta,
+                     vecY, compute_type, alg, rocsparse_spmv_stage_compute,
+                     &buffer_size, tmp_buffer));
+#elif KOKKOSSPARSE_IMPL_ROCM_VERSION >= 50400
   KOKKOS_ROCSPARSE_SAFE_CALL_IMPL(rocsparse_spmv_ex(
       handle, myRocsparseOperation, &alpha, Aspmat, vecX, &beta, vecY,
       compute_type, alg, rocsparse_spmv_stage_auto, &buffer_size, tmp_buffer));