diff --git a/CHANGELOG.md b/CHANGELOG.md index 6bfc00b4fe..3ebb102517 100644 --- 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 index 8e990cece5..4847b51e9b 100644 --- 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 index 49ae14ad9d..555c942c12 100644 --- 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, \ Kokkos::MemoryTraits>, \ true, ETI_SPEC_AVAIL> { \ - using XVector = Kokkos::View, LAYOUT, \ + using XVector = Kokkos::View*, LAYOUT, \ Kokkos::Device, \ Kokkos::MemoryTraits>; \ - using YVector = Kokkos::View, LAYOUT, \ + using YVector = Kokkos::View*, LAYOUT, \ Kokkos::Device, \ Kokkos::MemoryTraits>; \ static void swap(EXECSPACE const& space, XVector const& X, \ diff --git a/lapack/CMakeLists.txt b/lapack/CMakeLists.txt index 8ab784a325..7c0c3183bd 100644 --- 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 index b59d6d99c8..2ce9f69954 100644 --- 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 - -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 +#include namespace KokkosLapack { namespace Impl { diff --git a/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp b/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp index 2baa76a132..5846e177d6 100644 --- 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 +#include namespace KokkosLapack { namespace Impl { diff --git a/lapack/tpls/KokkosLapack_magma.hpp b/lapack/tpls/KokkosLapack_magma.hpp new file mode 100644 index 0000000000..66529d73de --- /dev/null +++ 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 index 3ed0623018..655b5b8579 100644 --- 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 index 2ece96fed6..26f95694e9 100644 --- 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 index cc34e55093..baf2d3a822 100644 --- a/sparse/src/KokkosSparse_Utils_rocsparse.hpp +++ b/sparse/src/KokkosSparse_Utils_rocsparse.hpp @@ -21,8 +21,12 @@ #include #ifdef KOKKOSKERNELS_ENABLE_TPL_ROCSPARSE +#if __has_include() +#include +#else #include -#include "rocsparse/rocsparse.h" +#endif +#include 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 index 97019e4682..75752190e7 100644 --- 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) { + 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) { + 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>) { + 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>) { + 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, + "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) { 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 index efb591375b..a4c50e437f 100644 --- 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, int64_t, size_t, // rocSPARSE #if defined(KOKKOSKERNELS_ENABLE_TPL_ROCSPARSE) -#include -#include #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));