From 5bd8705c78022b99342181facb24f51d6310f0b6 Mon Sep 17 00:00:00 2001 From: Nathan Ellingwood Date: Thu, 25 Apr 2024 11:44:08 -0600 Subject: [PATCH 01/19] Update to version 4.3.1 --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index bd3d761bdb..e575263e55 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 3) -SET(KokkosKernels_VERSION_PATCH 0) +SET(KokkosKernels_VERSION_PATCH 1) SET(KokkosKernels_VERSION "${KokkosKernels_VERSION_MAJOR}.${KokkosKernels_VERSION_MINOR}.${KokkosKernels_VERSION_PATCH}") #Set variables for config file From 8da022d00550719814031ac6cc351002ae353526 Mon Sep 17 00:00:00 2001 From: Nathan Ellingwood Date: Thu, 25 Apr 2024 11:48:04 -0600 Subject: [PATCH 02/19] CHANGELOG: fixups from develop (remove trailing white space) --- CHANGELOG.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 6bc9cb65a6..4e6da70740 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -639,7 +639,7 @@ ## [3.6.00](https://github.com/kokkos/kokkos-kernels/tree/3.6.00) (2022-02-18) [Full Changelog](https://github.com/kokkos/kokkos-kernels/compare/3.5.00...3.6.00) -### Features: +### Features: #### Batched Sparse Linear algebra - Kokkos Kernels is adding a new component to the library: batched sparse linear algebra. @@ -673,7 +673,7 @@ - SpMV: adding support for rocSPARSE TPL [\#1221](https://github.com/kokkos/kokkos-kernels/pull/1221) #### Additional new features -- bhalf: Unit test Batched GEMM [\#1251](https://github.com/kokkos/kokkos-kernels/pull/1251) +- bhalf: Unit test Batched GEMM [\#1251](https://github.com/kokkos/kokkos-kernels/pull/1251) - and demostrate GMRES example convergence with bhalf_t (https://github.com/kokkos/kokkos-kernels/pull/1300) - Stream interface: adding stream support in GEMV and GEMM [\#1131](https://github.com/kokkos/kokkos-kernels/pull/1131) - Improve double buffering batched gemm performance [\#1217](https://github.com/kokkos/kokkos-kernels/pull/1217) @@ -962,7 +962,7 @@ ## [3.1.01](https://github.com/kokkos/kokkos-kernels/tree/3.1.01) (2020-05-04) [Full Changelog](https://github.com/kokkos/kokkos-kernels/compare/3.1.00...3.1.01) -** Fixed bugs:** +** Fixed bugs:** - KokkosBatched QR PR breaking nightly tests [\#691](https://github.com/kokkos/kokkos-kernels/issues/691) From 40bae22c7d9e7bc915cc5eed04fd1cd5e19c5743 Mon Sep 17 00:00:00 2001 From: James Foucar Date: Tue, 9 Apr 2024 14:10:35 -0600 Subject: [PATCH 03/19] Jgfouca/block spiluk fixes (#2172) * Progress * Attempt 1, fix multiplication order * Converges in 1 step * Various cleanups * Be sure not to reduce performance of unblocked impl Also add some comments. * Remove test mangling * Fixes for GPU * Fix warning * formatting * Increase eps for floats * This is no longer needed --- .../impl/KokkosBatched_Trsm_Team_Impl.hpp | 46 +++- .../impl/KokkosSparse_spiluk_numeric_impl.hpp | 260 ++++++++++++++---- sparse/unit_test/Test_Sparse_spiluk.hpp | 12 +- 3 files changed, 248 insertions(+), 70 deletions(-) diff --git a/batched/dense/impl/KokkosBatched_Trsm_Team_Impl.hpp b/batched/dense/impl/KokkosBatched_Trsm_Team_Impl.hpp index a7430775ea..9f5f857e44 100644 --- a/batched/dense/impl/KokkosBatched_Trsm_Team_Impl.hpp +++ b/batched/dense/impl/KokkosBatched_Trsm_Team_Impl.hpp @@ -99,6 +99,48 @@ struct TeamTrsm +struct TeamTrsm { + template + KOKKOS_INLINE_FUNCTION static int invoke(const MemberType &member, + const ScalarType alpha, + const AViewType &A, + const BViewType &B) { + return TeamTrsmInternalLeftUpper::invoke( + member, ArgDiag::use_unit_diag, B.extent(1), B.extent(0), alpha, + A.data(), A.stride_1(), A.stride_0(), B.data(), B.stride_1(), + B.stride_0()); + } +}; + +template +struct TeamTrsm { + template + KOKKOS_INLINE_FUNCTION static int invoke(const MemberType &member, + const ScalarType alpha, + const AViewType &A, + const BViewType &B) { + return TeamTrsmInternalLeftUpper::invoke( + member, ArgDiag::use_unit_diag, B.extent(1), B.extent(0), alpha, + A.data(), A.stride_1(), A.stride_0(), B.data(), B.stride_1(), + B.stride_0()); + } +}; + +/// +/// R/U/T +/// +/// B := (alpha*B) inv(triu(A)) +/// A(n x n), B(m x n) + template struct TeamTrsm { @@ -107,7 +149,7 @@ struct TeamTrsm::invoke( + return TeamTrsmInternalLeftUpper::invoke( member, ArgDiag::use_unit_diag, B.extent(1), B.extent(0), alpha, A.data(), A.stride_0(), A.stride_1(), B.data(), B.stride_1(), B.stride_0()); @@ -122,7 +164,7 @@ struct TeamTrsm::invoke( + return TeamTrsmInternalLeftUpper::invoke( member, ArgDiag::use_unit_diag, B.extent(1), B.extent(0), alpha, A.data(), A.stride_0(), A.stride_1(), B.data(), B.stride_1(), B.stride_0()); diff --git a/sparse/impl/KokkosSparse_spiluk_numeric_impl.hpp b/sparse/impl/KokkosSparse_spiluk_numeric_impl.hpp index b3b5dfa277..415ccf87a0 100644 --- a/sparse/impl/KokkosSparse_spiluk_numeric_impl.hpp +++ b/sparse/impl/KokkosSparse_spiluk_numeric_impl.hpp @@ -32,6 +32,9 @@ #include "KokkosBatched_Gemm_Decl.hpp" #include "KokkosBatched_Gemm_Serial_Impl.hpp" #include "KokkosBlas1_set.hpp" +#include "KokkosBatched_LU_Decl.hpp" +#include "KokkosBatched_Trmm_Decl.hpp" +#include "KokkosBatched_Trmm_Serial_Impl.hpp" //#define NUMERIC_OUTPUT_INFO @@ -107,6 +110,17 @@ struct IlukWrap { lno_t lev_start; using reftype = scalar_t &; + using valtype = scalar_t; + + static constexpr size_type BUFF_SIZE = 1; + + struct SBlock { + template + KOKKOS_INLINE_FUNCTION SBlock(T, size_type, size_type) {} + + KOKKOS_INLINE_FUNCTION + scalar_t *data() { return nullptr; } + }; Common(const ARowMapType &A_row_map_, const AEntriesType &A_entries_, const AValuesType &A_values_, const LRowMapType &L_row_map_, @@ -131,6 +145,9 @@ struct IlukWrap { "Tried to use blocks with the unblocked Common?"); } + KOKKOS_INLINE_FUNCTION + size_type get_block_size() const { return 0; } + // lset KOKKOS_INLINE_FUNCTION void lset(const size_type nnz, const scalar_t &value) const { @@ -154,12 +171,18 @@ struct IlukWrap { // divide. lhs /= rhs KOKKOS_INLINE_FUNCTION - void divide(const member_type &team, scalar_t &lhs, - const scalar_t &rhs) const { + void divide(const member_type &team, scalar_t &lhs, const scalar_t &rhs, + scalar_t *) const { Kokkos::single(Kokkos::PerTeam(team), [&]() { lhs /= rhs; }); team.team_barrier(); } + // divide_left. lhs /= rhs + KOKKOS_INLINE_FUNCTION + void divide_left(scalar_t &lhs, const scalar_t &rhs, scalar_t *) const { + lhs /= rhs; + } + // multiply_subtract. C -= A * B KOKKOS_INLINE_FUNCTION void multiply_subtract(const scalar_t &A, const scalar_t &B, @@ -171,6 +194,18 @@ struct IlukWrap { KOKKOS_INLINE_FUNCTION scalar_t &lget(const size_type nnz) const { return L_values(nnz); } + // lcopy + KOKKOS_INLINE_FUNCTION + scalar_t lcopy(const size_type nnz, scalar_t *) const { + return L_values(nnz); + } + + // ucopy + KOKKOS_INLINE_FUNCTION + scalar_t ucopy(const size_type nnz, scalar_t *) const { + return U_values(nnz); + } + // uget KOKKOS_INLINE_FUNCTION scalar_t &uget(const size_type nnz) const { return U_values(nnz); } @@ -188,6 +223,12 @@ struct IlukWrap { // print KOKKOS_INLINE_FUNCTION void print(const scalar_t &item) const { std::cout << item << std::endl; } + + // report + KOKKOS_INLINE_FUNCTION + void report() const { + std::cout << "JGF using unblocked version" << std::endl; + } }; // Partial specialization for block support @@ -197,6 +238,30 @@ struct IlukWrap { struct Common { + // BSR data is in LayoutRight! + using Layout = Kokkos::LayoutRight; + using value_type = typename LValuesType::value_type; + using cvalue_type = typename LValuesType::const_value_type; + + using Block = Kokkos::View< + value_type **, Layout, typename LValuesType::device_type, + Kokkos::MemoryTraits >; + + // const block + using CBlock = Kokkos::View< + cvalue_type **, Layout, typename UValuesType::device_type, + Kokkos::MemoryTraits >; + + // scratch block + using SBlock = Kokkos::View< + value_type **, Layout, typename execution_space::scratch_memory_space, + Kokkos::MemoryTraits >; + + using reftype = Block; + using valtype = Block; + + static constexpr size_type BUFF_SIZE = 128; + ARowMapType A_row_map; AEntriesType A_entries; AValuesType A_values; @@ -212,26 +277,6 @@ struct IlukWrap { size_type block_size; size_type block_items; - // BSR data is in LayoutRight! - using Layout = Kokkos::LayoutRight; - - using LBlock = Kokkos::View< - typename LValuesType::value_type **, Layout, - typename LValuesType::device_type, - Kokkos::MemoryTraits >; - - using UBlock = Kokkos::View< - typename UValuesType::value_type **, Layout, - typename UValuesType::device_type, - Kokkos::MemoryTraits >; - - using ABlock = Kokkos::View< - typename AValuesType::value_type **, Layout, - typename AValuesType::device_type, - Kokkos::MemoryTraits >; - - using reftype = LBlock; - Common(const ARowMapType &A_row_map_, const AEntriesType &A_entries_, const AValuesType &A_values_, const LRowMapType &L_row_map_, const LEntriesType &L_entries_, LValuesType &L_values_, @@ -255,8 +300,12 @@ struct IlukWrap { block_items(block_size * block_size) { KK_REQUIRE_MSG(block_size > 0, "Tried to use block_size=0 with the blocked Common?"); + KK_REQUIRE_MSG(block_size <= 11, "Max supported block size is 11"); } + KOKKOS_INLINE_FUNCTION + size_type get_block_size() const { return block_size; } + // lset KOKKOS_INLINE_FUNCTION void lset(const size_type block, const scalar_t &value) const { @@ -264,13 +313,9 @@ struct IlukWrap { } KOKKOS_INLINE_FUNCTION - void lset(const size_type block, const ABlock &rhs) const { + void lset(const size_type block, const CBlock &rhs) const { auto lblock = lget(block); - for (size_type i = 0; i < block_size; ++i) { - for (size_type j = 0; j < block_size; ++j) { - lblock(i, j) = rhs(i, j); - } - } + assign(lblock, rhs); } // uset @@ -280,13 +325,9 @@ struct IlukWrap { } KOKKOS_INLINE_FUNCTION - void uset(const size_type block, const ABlock &rhs) const { + void uset(const size_type block, const CBlock &rhs) const { auto ublock = uget(block); - for (size_type i = 0; i < block_size; ++i) { - for (size_type j = 0; j < block_size; ++j) { - ublock(i, j) = rhs(i, j); - } - } + assign(ublock, rhs); } // lset_id @@ -295,49 +336,111 @@ struct IlukWrap { KokkosBatched::TeamSetIdentity::invoke(team, lget(block)); } - // divide. lhs /= rhs + // assign + template + KOKKOS_INLINE_FUNCTION void assign(const ViewT &lhs, + const CBlock &rhs) const { + for (size_type i = 0; i < block_size; ++i) { + for (size_type j = 0; j < block_size; ++j) { + lhs(i, j) = rhs(i, j); + } + } + } + + // divide. lhs /= rhs (lhs = lhs * rhs^-1) KOKKOS_INLINE_FUNCTION - void divide(const member_type &team, const LBlock &lhs, - const UBlock &rhs) const { + void divide(const member_type &team, const Block &lhs, const CBlock &rhs, + scalar_t *buff) const { + // Need a temp block to do LU of rhs + Block LU(buff, block_size, block_size); + assign(LU, rhs); + KokkosBatched::TeamLU::invoke(team, LU); + + // rhs = LU + // rhs^-1 = U^-1 * L^-1 + // lhs = (lhs * U^-1) * L^-1, so do U trsm first KokkosBatched::TeamTrsm< member_type, KokkosBatched::Side::Right, KokkosBatched::Uplo::Upper, - KokkosBatched::Trans::NoTranspose, // not 100% on this - KokkosBatched::Diag::NonUnit, - KokkosBatched::Algo::Trsm::Unblocked>:: // not 100% on this - invoke(team, 1.0, rhs, lhs); + KokkosBatched::Trans::NoTranspose, KokkosBatched::Diag::NonUnit, + KokkosBatched::Algo::Trsm::Blocked>::invoke(team, 1.0, LU, lhs); + + KokkosBatched::TeamTrsm< + member_type, KokkosBatched::Side::Right, KokkosBatched::Uplo::Lower, + KokkosBatched::Trans::NoTranspose, KokkosBatched::Diag::Unit, + KokkosBatched::Algo::Trsm::Blocked>::invoke(team, 1.0, LU, lhs); + } + + // divide_left. lhs /= rhs (lhs = rhs^-1 * lhs) + KOKKOS_INLINE_FUNCTION + void divide_left(const Block &lhs, const CBlock &rhs, + scalar_t *buff) const { + Block LU(buff, block_size, block_size); + assign(LU, rhs); + KokkosBatched::SerialLU::invoke(LU); + + // rhs = LU + // rhs^-1 = U^-1 * L^-1 + // lhs = U^-1 * (L^-1 * lhs), so do L trsm first + KokkosBatched::SerialTrsm< + KokkosBatched::Side::Left, KokkosBatched::Uplo::Lower, + KokkosBatched::Trans::NoTranspose, KokkosBatched::Diag::Unit, + KokkosBatched::Algo::Trsm::Blocked>::invoke(1.0, LU, lhs); + + KokkosBatched::SerialTrsm< + KokkosBatched::Side::Left, KokkosBatched::Uplo::Upper, + KokkosBatched::Trans::NoTranspose, KokkosBatched::Diag::NonUnit, + KokkosBatched::Algo::Trsm::Blocked>::invoke(1.0, LU, lhs); } // multiply_subtract. C -= A * B - template - KOKKOS_INLINE_FUNCTION void multiply_subtract(const UBlock &A, - const LBlock &B, - CView &C) const { + KOKKOS_INLINE_FUNCTION + void multiply_subtract(const CBlock &A, const CBlock &B, + const Block &C) const { // Use gemm. alpha is hardcoded to -1, beta hardcoded to 1 KokkosBatched::SerialGemm< KokkosBatched::Trans::NoTranspose, KokkosBatched::Trans::NoTranspose, - KokkosBatched::Algo::Gemm::Unblocked>::invoke( - -1.0, A, B, 1.0, C); + KokkosBatched::Algo::Gemm::Blocked>::invoke(-1.0, A, B, 1.0, + C); } // lget KOKKOS_INLINE_FUNCTION - LBlock lget(const size_type block) const { - return LBlock(L_values.data() + (block * block_items), block_size, - block_size); + Block lget(const size_type block) const { + return Block(L_values.data() + (block * block_items), block_size, + block_size); + } + + // lcopy + KOKKOS_INLINE_FUNCTION + Block lcopy(const size_type block, scalar_t *buff) const { + Block result(buff, block_size, block_size); + auto lblock = lget(block); + assign(result, lblock); + return result; + } + + // ucopy + KOKKOS_INLINE_FUNCTION + Block ucopy(const size_type block, scalar_t *buff) const { + Block result(buff, block_size, block_size); + auto ublock = uget(block); + assign(result, ublock); + return result; } // uget KOKKOS_INLINE_FUNCTION - UBlock uget(const size_type block) const { - return UBlock(U_values.data() + (block * block_items), block_size, - block_size); + Block uget(const size_type block) const { + return Block(U_values.data() + (block * block_items), block_size, + block_size); } // aget KOKKOS_INLINE_FUNCTION - ABlock aget(const size_type block) const { - return ABlock(A_values.data() + (block * block_items), block_size, + CBlock aget(const size_type block) const { + return CBlock(A_values.data() + (block * block_items), block_size, block_size); } @@ -357,7 +460,7 @@ struct IlukWrap { // print KOKKOS_INLINE_FUNCTION - void print(const LBlock &item) const { + void print(const CBlock &item) const { for (size_type i = 0; i < block_size; ++i) { std::cout << " "; for (size_type j = 0; j < block_size; ++j) { @@ -366,6 +469,13 @@ struct IlukWrap { std::cout << std::endl; } } + + // report + KOKKOS_INLINE_FUNCTION + void report() const { + std::cout << "JGF using blocked version with block_size=" << block_size + << std::endl; + } }; template struct SpilukTest { @@ -130,6 +128,7 @@ struct SpilukTest { using EntriesType = Kokkos::View; using ValuesType = Kokkos::View; using AT = Kokkos::ArithTraits; + using mag_t = typename Kokkos::ArithTraits::mag_type; using RowMapType_hostmirror = typename RowMapType::HostMirror; using EntriesType_hostmirror = typename EntriesType::HostMirror; @@ -138,6 +137,9 @@ struct SpilukTest { using memory_space = typename device::memory_space; using range_policy = Kokkos::RangePolicy; + static constexpr double EPS = + std::is_same::value ? 1e-7 : 1e-4; + using KernelHandle = KokkosKernels::Experimental::KokkosKernelsHandle< size_type, lno_t, scalar_t, execution_space, memory_space, memory_space>; @@ -243,11 +245,7 @@ struct SpilukTest { } if (fill_lev > 1) { - if (UseBlocks) { - EXPECT_LT(result, 1e-2); - } else { - EXPECT_LT(result, 1e-4); - } + EXPECT_LT(result, 1e-4); } } From 2544c9d9be2e256987f7c284811ce64b612faef2 Mon Sep 17 00:00:00 2001 From: Nathan Ellingwood Date: Fri, 12 Apr 2024 16:33:56 -0600 Subject: [PATCH 04/19] Merge pull request #2176 from ndellingwood/issue-2175 Add guard for cusparse spmv_mv_tpl_spec_avail (cherry picked from commit cb824ca0cdf32fe6508b5ebe092d9e90cdaa6cfd) --- sparse/tpls/KokkosSparse_spmv_mv_tpl_spec_avail.hpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sparse/tpls/KokkosSparse_spmv_mv_tpl_spec_avail.hpp b/sparse/tpls/KokkosSparse_spmv_mv_tpl_spec_avail.hpp index 88fef4421a..44a8098ca3 100644 --- a/sparse/tpls/KokkosSparse_spmv_mv_tpl_spec_avail.hpp +++ b/sparse/tpls/KokkosSparse_spmv_mv_tpl_spec_avail.hpp @@ -29,6 +29,7 @@ struct spmv_mv_tpl_spec_avail { enum : bool { value = false }; }; +#ifdef KOKKOSKERNELS_ENABLE_TPL_CUSPARSE #define KOKKOSSPARSE_SPMV_MV_TPL_SPEC_AVAIL_CUSPARSE(SCALAR, ORDINAL, OFFSET, \ XL, YL, MEMSPACE) \ template <> \ @@ -152,6 +153,7 @@ KOKKOSSPARSE_SPMV_MV_TPL_SPEC_AVAIL_CUSPARSE(Kokkos::Experimental::half_t, int, #endif #endif // defined(CUSPARSE_VERSION) && (10300 <= CUSPARSE_VERSION) +#endif } // namespace Impl } // namespace KokkosSparse From 1cd3d5099200106d08a7dc1f9f4ca45c21687637 Mon Sep 17 00:00:00 2001 From: Nathan Ellingwood Date: Thu, 18 Apr 2024 11:07:15 -0600 Subject: [PATCH 05/19] Resolves multiple definition of Magma and Cuda singletons (#2178) Address issue #2175 --- blas/tpls/KokkosBlas_Cuda_tpl.cpp | 1 + blas/tpls/KokkosBlas_Cuda_tpl.hpp | 22 ---------- blas/tpls/KokkosBlas_Magma_tpl.hpp | 41 +++++++++++++++++++ blas/tpls/KokkosBlas_magma.hpp | 37 +++++++++++++++++ blas/tpls/KokkosBlas_tpl_spec.hpp | 17 -------- lapack/tpls/KokkosLapack_Cuda_tpl.cpp | 1 + lapack/tpls/KokkosLapack_Cuda_tpl.hpp | 22 ---------- lapack/tpls/KokkosLapack_Magma_tpl.hpp | 41 +++++++++++++++++++ .../tpls/KokkosLapack_gesv_tpl_spec_decl.hpp | 2 +- 9 files changed, 122 insertions(+), 62 deletions(-) create mode 100644 blas/tpls/KokkosBlas_Magma_tpl.hpp create mode 100644 blas/tpls/KokkosBlas_magma.hpp create mode 100644 lapack/tpls/KokkosLapack_Magma_tpl.hpp diff --git a/blas/tpls/KokkosBlas_Cuda_tpl.cpp b/blas/tpls/KokkosBlas_Cuda_tpl.cpp index eed90ef7e0..cb8ba34101 100644 --- a/blas/tpls/KokkosBlas_Cuda_tpl.cpp +++ b/blas/tpls/KokkosBlas_Cuda_tpl.cpp @@ -16,3 +16,4 @@ #include #include #include +#include diff --git a/blas/tpls/KokkosBlas_Cuda_tpl.hpp b/blas/tpls/KokkosBlas_Cuda_tpl.hpp index cf51341471..d85785316e 100644 --- a/blas/tpls/KokkosBlas_Cuda_tpl.hpp +++ b/blas/tpls/KokkosBlas_Cuda_tpl.hpp @@ -39,26 +39,4 @@ CudaBlasSingleton& CudaBlasSingleton::singleton() { } // namespace KokkosBlas #endif // defined (KOKKOSKERNELS_ENABLE_TPL_CUBLAS) -#if defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) -#include - -namespace KokkosBlas { -namespace Impl { - -MagmaSingleton::MagmaSingleton() { - magma_int_t stat = magma_init(); - if (stat != MAGMA_SUCCESS) Kokkos::abort("MAGMA initialization failed\n"); - - Kokkos::push_finalize_hook([&]() { magma_finalize(); }); -} - -MagmaSingleton& MagmaSingleton::singleton() { - static MagmaSingleton s; - return s; -} - -} // namespace Impl -} // namespace KokkosBlas -#endif // defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) - #endif // KOKKOSBLAS_CUDA_TPL_HPP_ diff --git a/blas/tpls/KokkosBlas_Magma_tpl.hpp b/blas/tpls/KokkosBlas_Magma_tpl.hpp new file mode 100644 index 0000000000..f149a790df --- /dev/null +++ b/blas/tpls/KokkosBlas_Magma_tpl.hpp @@ -0,0 +1,41 @@ +//@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 KOKKOSBLAS_MAGMA_TPL_HPP_ +#define KOKKOSBLAS_MAGMA_TPL_HPP_ + +#if defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) +#include + +namespace KokkosBlas { +namespace Impl { + +MagmaSingleton::MagmaSingleton() { + magma_int_t stat = magma_init(); + if (stat != MAGMA_SUCCESS) Kokkos::abort("MAGMA initialization failed\n"); + + Kokkos::push_finalize_hook([&]() { magma_finalize(); }); +} + +MagmaSingleton& MagmaSingleton::singleton() { + static MagmaSingleton s; + return s; +} + +} // namespace Impl +} // namespace KokkosBlas +#endif // defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) + +#endif // KOKKOSBLAS_MAGMA_TPL_HPP_ diff --git a/blas/tpls/KokkosBlas_magma.hpp b/blas/tpls/KokkosBlas_magma.hpp new file mode 100644 index 0000000000..5f5fcfe4e1 --- /dev/null +++ b/blas/tpls/KokkosBlas_magma.hpp @@ -0,0 +1,37 @@ +//@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 KOKKOSBLAS_MAGMA_HPP_ +#define KOKKOSBLAS_MAGMA_HPP_ + +// If LAPACK TPL is enabled, it is preferred over magma's LAPACK +#ifdef KOKKOSKERNELS_ENABLE_TPL_MAGMA +#include "magma_v2.h" + +namespace KokkosBlas { +namespace Impl { + +struct MagmaSingleton { + MagmaSingleton(); + + static MagmaSingleton& singleton(); +}; + +} // namespace Impl +} // namespace KokkosBlas +#endif // KOKKOSKERNELS_ENABLE_TPL_MAGMA + +#endif // KOKKOSBLAS_MAGMA_HPP_ diff --git a/blas/tpls/KokkosBlas_tpl_spec.hpp b/blas/tpls/KokkosBlas_tpl_spec.hpp index a1eee4b69c..0151c0534f 100644 --- a/blas/tpls/KokkosBlas_tpl_spec.hpp +++ b/blas/tpls/KokkosBlas_tpl_spec.hpp @@ -214,21 +214,4 @@ inline rocblas_operation trans_mode_kk_to_rocblas(const char kkMode[]) { #endif // KOKKOSKERNELS_ENABLE_TPL_ROCBLAS -// If LAPACK TPL is enabled, it is preferred over magma's LAPACK -#ifdef KOKKOSKERNELS_ENABLE_TPL_MAGMA -#include "magma_v2.h" - -namespace KokkosBlas { -namespace Impl { - -struct MagmaSingleton { - MagmaSingleton(); - - static MagmaSingleton& singleton(); -}; - -} // namespace Impl -} // namespace KokkosBlas -#endif // KOKKOSKERNELS_ENABLE_TPL_MAGMA - #endif // KOKKOSBLAS_TPL_SPEC_HPP_ diff --git a/lapack/tpls/KokkosLapack_Cuda_tpl.cpp b/lapack/tpls/KokkosLapack_Cuda_tpl.cpp index 2ac28871a4..7f87eed0d5 100644 --- a/lapack/tpls/KokkosLapack_Cuda_tpl.cpp +++ b/lapack/tpls/KokkosLapack_Cuda_tpl.cpp @@ -16,3 +16,4 @@ #include #include #include +#include diff --git a/lapack/tpls/KokkosLapack_Cuda_tpl.hpp b/lapack/tpls/KokkosLapack_Cuda_tpl.hpp index 6749a4740f..943d10d111 100644 --- a/lapack/tpls/KokkosLapack_Cuda_tpl.hpp +++ b/lapack/tpls/KokkosLapack_Cuda_tpl.hpp @@ -39,26 +39,4 @@ CudaLapackSingleton& CudaLapackSingleton::singleton() { } // namespace KokkosLapack #endif // defined (KOKKOSKERNELS_ENABLE_TPL_CUSOLVER) -#if defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) -#include - -namespace KokkosLapack { -namespace Impl { - -MagmaSingleton::MagmaSingleton() { - magma_int_t stat = magma_init(); - if (stat != MAGMA_SUCCESS) Kokkos::abort("MAGMA initialization failed\n"); - - Kokkos::push_finalize_hook([&]() { magma_finalize(); }); -} - -MagmaSingleton& MagmaSingleton::singleton() { - static MagmaSingleton s; - return s; -} - -} // namespace Impl -} // namespace KokkosLapack -#endif // defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) - #endif // KOKKOSLAPACK_CUDA_TPL_HPP_ diff --git a/lapack/tpls/KokkosLapack_Magma_tpl.hpp b/lapack/tpls/KokkosLapack_Magma_tpl.hpp new file mode 100644 index 0000000000..636c40735d --- /dev/null +++ b/lapack/tpls/KokkosLapack_Magma_tpl.hpp @@ -0,0 +1,41 @@ +//@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_TPL_HPP_ +#define KOKKOSLAPACK_MAGMA_TPL_HPP_ + +#if defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) +#include + +namespace KokkosLapack { +namespace Impl { + +MagmaSingleton::MagmaSingleton() { + magma_int_t stat = magma_init(); + if (stat != MAGMA_SUCCESS) Kokkos::abort("MAGMA initialization failed\n"); + + Kokkos::push_finalize_hook([&]() { magma_finalize(); }); +} + +MagmaSingleton& MagmaSingleton::singleton() { + static MagmaSingleton s; + return s; +} + +} // namespace Impl +} // namespace KokkosLapack +#endif // defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) + +#endif // KOKKOSLAPACK_MAGMA_TPL_HPP_ diff --git a/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp b/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp index 41592e079a..ca4b9e7abc 100644 --- a/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp +++ b/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp @@ -155,7 +155,7 @@ KOKKOSLAPACK_GESV_LAPACK(Kokkos::complex, Kokkos::LayoutLeft, // MAGMA #ifdef KOKKOSKERNELS_ENABLE_TPL_MAGMA -#include +#include namespace KokkosLapack { namespace Impl { From e91f892e5d8559c8775a1349ce7766214daf7460 Mon Sep 17 00:00:00 2001 From: Nathan Ellingwood Date: Thu, 18 Apr 2024 16:01:37 -0600 Subject: [PATCH 06/19] Merge pull request #2180 from MalachiTimothyPhillips/malachi/fix-divide-by-zero-in-trsv Add early return if numRows == 0 in trsv to avoid integer divide-by-zero error (cherry picked from commit 83374bfde15b549068f684ba9542e8fa09896473) --- sparse/impl/KokkosSparse_trsv_impl.hpp | 48 +++++++++++++++++++------- 1 file changed, 36 insertions(+), 12 deletions(-) diff --git a/sparse/impl/KokkosSparse_trsv_impl.hpp b/sparse/impl/KokkosSparse_trsv_impl.hpp index 9adb029d12..6d04a42877 100644 --- a/sparse/impl/KokkosSparse_trsv_impl.hpp +++ b/sparse/impl/KokkosSparse_trsv_impl.hpp @@ -184,7 +184,9 @@ struct TrsvWrap { static void lowerTriSolveCsrUnitDiag(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; const lno_t numVecs = X.extent(1); @@ -211,7 +213,9 @@ struct TrsvWrap { static void lowerTriSolveCsr(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; const lno_t numVecs = X.extent(1); @@ -254,7 +258,9 @@ struct TrsvWrap { static void upperTriSolveCsrUnitDiag(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; const lno_t numVecs = X.extent(1); @@ -304,7 +310,9 @@ struct TrsvWrap { static void upperTriSolveCsr(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; const lno_t numVecs = X.extent(1); @@ -371,7 +379,9 @@ struct TrsvWrap { static void upperTriSolveCscUnitDiag(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numCols = A.numCols(); const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; @@ -422,7 +432,9 @@ struct TrsvWrap { static void upperTriSolveCsc(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numCols = A.numCols(); const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; @@ -481,7 +493,9 @@ struct TrsvWrap { static void lowerTriSolveCscUnitDiag(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numCols = A.numCols(); const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; @@ -510,7 +524,9 @@ struct TrsvWrap { static void upperTriSolveCscUnitDiagConj(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numCols = A.numCols(); const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; @@ -562,7 +578,9 @@ struct TrsvWrap { static void upperTriSolveCscConj(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numCols = A.numCols(); const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; @@ -620,7 +638,9 @@ struct TrsvWrap { static void lowerTriSolveCsc(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numCols = A.numCols(); const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; @@ -657,7 +677,9 @@ struct TrsvWrap { static void lowerTriSolveCscUnitDiagConj(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numCols = A.numCols(); const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; @@ -686,7 +708,9 @@ struct TrsvWrap { static void lowerTriSolveCscConj(RangeMultiVectorType X, const CrsMatrixType& A, DomainMultiVectorType Y) { - const lno_t numRows = A.numRows(); + const lno_t numRows = A.numRows(); + if (numRows == 0) return; + const lno_t numCols = A.numCols(); const lno_t numPointRows = A.numPointRows(); const lno_t block_size = numPointRows / numRows; From b8f475ea272307ea85043332e645ebd3e37c3736 Mon Sep 17 00:00:00 2001 From: Nathan Ellingwood Date: Tue, 23 Apr 2024 16:19:47 -0600 Subject: [PATCH 07/19] magma: fix linker errors for builds without cusolver (#2181) * magma: fix linker errors for builds without cusolver * BatchedGemm test: workaround testing cublas+magma - temporary workaround to skip magma test when cublas enabled to avoid issues like #2177 --- .../unit_test/Test_Batched_BatchedGemm.hpp | 6 +++++- lapack/CMakeLists.txt | 6 ++++++ lapack/tpls/KokkosLapack_Cuda_tpl.cpp | 1 - lapack/tpls/KokkosLapack_Magma_tpl.cpp | 18 ++++++++++++++++++ 4 files changed, 29 insertions(+), 2 deletions(-) create mode 100644 lapack/tpls/KokkosLapack_Magma_tpl.cpp diff --git a/batched/dense/unit_test/Test_Batched_BatchedGemm.hpp b/batched/dense/unit_test/Test_Batched_BatchedGemm.hpp index d57e671908..3c00b4f477 100644 --- a/batched/dense/unit_test/Test_Batched_BatchedGemm.hpp +++ b/batched/dense/unit_test/Test_Batched_BatchedGemm.hpp @@ -229,7 +229,11 @@ void impl_test_batched_gemm(const int N, const int matAdim1, const int matAdim2, ASSERT_EQ(batchedGemmHandleCublas.vecLen, 0); #endif -#if defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) + // FIXME temporary workaround to run this magma test only if cublas is not + // enabled the design of the BatchedGemmHandle currently does not allow + // simultanous testing in this way. See issue #2177 +#if defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) && \ + !defined(KOKKOSKERNELS_ENABLE_TPL_CUBLAS) magma_queue_t magma_queue; BatchedGemmHandle batchedGemmHandleMagma(magma_queue, GemmTplAlgos::MAGMA, 0, 0); diff --git a/lapack/CMakeLists.txt b/lapack/CMakeLists.txt index f825a2184a..804a2b7542 100644 --- a/lapack/CMakeLists.txt +++ b/lapack/CMakeLists.txt @@ -34,6 +34,12 @@ IF (KOKKOSKERNELS_ENABLE_TPL_CUSOLVER) ) ENDIF() +IF (KOKKOSKERNELS_ENABLE_TPL_MAGMA) + LIST(APPEND SOURCES + lapack/tpls/KokkosLapack_Magma_tpl.cpp + ) +ENDIF() + ################## # # # ETI generation # diff --git a/lapack/tpls/KokkosLapack_Cuda_tpl.cpp b/lapack/tpls/KokkosLapack_Cuda_tpl.cpp index 7f87eed0d5..2ac28871a4 100644 --- a/lapack/tpls/KokkosLapack_Cuda_tpl.cpp +++ b/lapack/tpls/KokkosLapack_Cuda_tpl.cpp @@ -16,4 +16,3 @@ #include #include #include -#include diff --git a/lapack/tpls/KokkosLapack_Magma_tpl.cpp b/lapack/tpls/KokkosLapack_Magma_tpl.cpp new file mode 100644 index 0000000000..73add8d9e0 --- /dev/null +++ b/lapack/tpls/KokkosLapack_Magma_tpl.cpp @@ -0,0 +1,18 @@ +//@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 +#include +#include +#include From be2a230cbf060e0ef211ec32875a13fc6779591a Mon Sep 17 00:00:00 2001 From: Malachi Date: Thu, 25 Apr 2024 09:57:59 -0500 Subject: [PATCH 08/19] Resolve vortex compilation issue by resolving (potentially) duplicate symbol (#2183) Stick to pattern of removing leading 'c' or 'z' in method name and relying on the template type Co-authored-by: malphil --- .../KokkosBlas2_syr2_tpl_spec_decl_blas.hpp | 4 ++-- .../KokkosBlas2_syr_tpl_spec_decl_blas.hpp | 4 ++-- blas/tpls/KokkosBlas_Host_tpl.cpp | 24 ++++++++++--------- blas/tpls/KokkosBlas_Host_tpl.hpp | 15 ++++-------- 4 files changed, 21 insertions(+), 26 deletions(-) diff --git a/blas/tpls/KokkosBlas2_syr2_tpl_spec_decl_blas.hpp b/blas/tpls/KokkosBlas2_syr2_tpl_spec_decl_blas.hpp index 8561675c72..f22e800bc5 100644 --- a/blas/tpls/KokkosBlas2_syr2_tpl_spec_decl_blas.hpp +++ b/blas/tpls/KokkosBlas2_syr2_tpl_spec_decl_blas.hpp @@ -163,7 +163,7 @@ namespace Impl { ETI_SPEC_AVAIL>::syr2(space, trans, uplo, alpha, X, Y, A); \ } else { \ if (A_is_ll) { \ - HostBlas>::zher2( \ + HostBlas>::her2( \ uplo[0], N, alpha, \ reinterpret_cast*>(X.data()), one, \ reinterpret_cast*>(Y.data()), one, \ @@ -220,7 +220,7 @@ namespace Impl { ETI_SPEC_AVAIL>::syr2(space, trans, uplo, alpha, X, Y, A); \ } else { \ if (A_is_ll) { \ - HostBlas>::cher2( \ + HostBlas>::her2( \ uplo[0], N, alpha, \ reinterpret_cast*>(X.data()), one, \ reinterpret_cast*>(Y.data()), one, \ diff --git a/blas/tpls/KokkosBlas2_syr_tpl_spec_decl_blas.hpp b/blas/tpls/KokkosBlas2_syr_tpl_spec_decl_blas.hpp index 6b64fce2bc..fc8fb949d7 100644 --- a/blas/tpls/KokkosBlas2_syr_tpl_spec_decl_blas.hpp +++ b/blas/tpls/KokkosBlas2_syr_tpl_spec_decl_blas.hpp @@ -139,7 +139,7 @@ namespace Impl { space, trans, uplo, alpha, X, A); \ } else { \ if (A_is_ll) { \ - HostBlas>::zher( \ + HostBlas>::her( \ uplo[0], N, alpha.real(), \ reinterpret_cast*>(X.data()), one, \ reinterpret_cast*>(A.data()), LDA); \ @@ -188,7 +188,7 @@ namespace Impl { space, trans, uplo, alpha, X, A); \ } else { \ if (A_is_ll && (alpha.imag() == 0.)) { \ - HostBlas>::cher( \ + HostBlas>::her( \ uplo[0], N, alpha.real(), \ reinterpret_cast*>(X.data()), one, \ reinterpret_cast*>(A.data()), LDA); \ diff --git a/blas/tpls/KokkosBlas_Host_tpl.cpp b/blas/tpls/KokkosBlas_Host_tpl.cpp index 50aab57c73..dc04ca7e67 100644 --- a/blas/tpls/KokkosBlas_Host_tpl.cpp +++ b/blas/tpls/KokkosBlas_Host_tpl.cpp @@ -295,10 +295,10 @@ void F77_BLAS_MANGLE(dsyr, DSYR)(const char*, KK_INT*, const double*, void F77_BLAS_MANGLE(cher, CHER)(const char*, KK_INT*, const float*, const std::complex*, KK_INT*, - std::complex*, KK_INT*); + /* */ std::complex*, KK_INT*); void F77_BLAS_MANGLE(zher, ZHER)(const char*, KK_INT*, const double*, const std::complex*, KK_INT*, - std::complex*, KK_INT*); + /* */ std::complex*, KK_INT*); /// /// Syr2 @@ -322,12 +322,12 @@ void F77_BLAS_MANGLE(cher2, CHER2)(const char*, KK_INT*, const std::complex*, const std::complex*, KK_INT*, const std::complex*, KK_INT*, - std::complex*, KK_INT*); + /* */ std::complex*, KK_INT*); void F77_BLAS_MANGLE(zher2, ZHER2)(const char*, KK_INT*, const std::complex*, const std::complex*, KK_INT*, const std::complex*, KK_INT*, - std::complex*, KK_INT*); + /* */ std::complex*, KK_INT*); /// /// Trsv @@ -901,14 +901,14 @@ void HostBlas >::gerc( } template <> template <> -void HostBlas >::cher( +void HostBlas >::her( const char uplo, KK_INT n, const float alpha, const std::complex* x, KK_INT incx, std::complex* a, KK_INT lda) { F77_FUNC_CHER(&uplo, &n, &alpha, (const std::complex*)x, &incx, (std::complex*)a, &lda); } template <> -void HostBlas >::cher2( +void HostBlas >::her2( const char uplo, KK_INT n, const std::complex alpha, const std::complex* x, KK_INT incx, const std::complex* y, KK_INT incy, std::complex* a, KK_INT lda) { @@ -1069,15 +1069,17 @@ void HostBlas >::gerc( } template <> template <> -void HostBlas >::zher( - const char uplo, KK_INT n, const double alpha, - const std::complex* x, KK_INT incx, std::complex* a, - KK_INT lda) { +void HostBlas >::her(const char uplo, KK_INT n, + const double alpha, + const std::complex* x, + KK_INT incx, + std::complex* a, + KK_INT lda) { F77_FUNC_ZHER(&uplo, &n, &alpha, (const std::complex*)x, &incx, (std::complex*)a, &lda); } template <> -void HostBlas >::zher2( +void HostBlas >::her2( const char uplo, KK_INT n, const std::complex alpha, const std::complex* x, KK_INT incx, const std::complex* y, KK_INT incy, std::complex* a, KK_INT lda) { diff --git a/blas/tpls/KokkosBlas_Host_tpl.hpp b/blas/tpls/KokkosBlas_Host_tpl.hpp index 5fb7c1f624..d28f7a2186 100644 --- a/blas/tpls/KokkosBlas_Host_tpl.hpp +++ b/blas/tpls/KokkosBlas_Host_tpl.hpp @@ -90,18 +90,11 @@ struct HostBlas { KK_INT incx, const T *y, KK_INT incy, T *a, KK_INT lda); template - static void cher(const char uplo, KK_INT n, const tAlpha alpha, const T *x, - KK_INT incx, T *a, KK_INT lda); - - template - static void zher(const char uplo, KK_INT n, const tAlpha alpha, const T *x, - KK_INT incx, T *a, KK_INT lda); - - static void cher2(const char uplo, KK_INT n, const T alpha, const T *x, - KK_INT incx, const T *y, KK_INT incy, T *a, KK_INT lda); + static void her(const char uplo, KK_INT n, const tAlpha alpha, const T *x, + KK_INT incx, T *a, KK_INT lda); - static void zher2(const char uplo, KK_INT n, const T alpha, const T *x, - KK_INT incx, const T *y, KK_INT incy, T *a, KK_INT lda); + static void her2(const char uplo, KK_INT n, const T alpha, const T *x, + KK_INT incx, const T *y, KK_INT incy, T *a, KK_INT lda); static void trsv(const char uplo, const char transa, const char diag, KK_INT m, const T *a, KK_INT lda, From 37bebb910259e7f5da5c5f8f63f23e91e9a39d2f Mon Sep 17 00:00:00 2001 From: brian-kelley Date: Wed, 3 Apr 2024 12:09:40 -0600 Subject: [PATCH 09/19] Fix #2156 (#2164) spmv: add special path for rank-2 x/y, but where both have 1 column and a TPL is available for rank-1 but not rank-2. Also call "subhandle->set_exec_space" correctly in the TPLs to ensure proper synchronization between setup, spmv and cleanup (in the case that different exec instances are used in different calls) --- sparse/src/KokkosSparse_spmv.hpp | 54 +++++++++++++++++++ sparse/src/KokkosSparse_spmv_handle.hpp | 5 +- ...kosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp | 5 ++ .../KokkosSparse_spmv_mv_tpl_spec_decl.hpp | 1 + .../tpls/KokkosSparse_spmv_tpl_spec_decl.hpp | 6 +++ 5 files changed, 68 insertions(+), 3 deletions(-) diff --git a/sparse/src/KokkosSparse_spmv.hpp b/sparse/src/KokkosSparse_spmv.hpp index 2391291695..f11b61f675 100644 --- a/sparse/src/KokkosSparse_spmv.hpp +++ b/sparse/src/KokkosSparse_spmv.hpp @@ -40,6 +40,31 @@ struct RANK_ONE {}; struct RANK_TWO {}; } // namespace +namespace Impl { +template +inline constexpr bool spmv_general_tpl_avail() { + constexpr bool isBSR = ::KokkosSparse::Experimental::is_bsr_matrix_v; + if constexpr (!isBSR) { + // CRS + if constexpr (XVector::rank() == 1) + return spmv_tpl_spec_avail::value; + else + return spmv_mv_tpl_spec_avail::value; + } else { + // BSR + if constexpr (XVector::rank() == 1) + return spmv_bsrmatrix_tpl_spec_avail::value; + else + return spmv_mv_bsrmatrix_tpl_spec_avail::value; + } +} +} // namespace Impl + // clang-format off /// \brief Kokkos sparse matrix-vector multiply. /// Computes y := alpha*Op(A)*x + beta*y, where Op(A) is @@ -221,6 +246,35 @@ void spmv(const ExecutionSpace& space, Handle* handle, const char mode[], typename KokkosKernels::Impl::GetUnifiedLayout::array_layout, typename YVector::device_type, Kokkos::MemoryTraits>; + // Special case: XVector/YVector are rank-2 but x,y both have one column and + // are contiguous. If a TPL is available for rank-1 vectors but not rank-2, + // take rank-1 subviews of x,y and call the rank-1 version. + if constexpr (XVector::rank() == 2) { + using XVector_SubInternal = Kokkos::View< + typename XVector::const_value_type*, + typename KokkosKernels::Impl::GetUnifiedLayout::array_layout, + typename XVector::device_type, + Kokkos::MemoryTraits>; + using YVector_SubInternal = Kokkos::View< + typename YVector::non_const_value_type*, + typename KokkosKernels::Impl::GetUnifiedLayout::array_layout, + typename YVector::device_type, Kokkos::MemoryTraits>; + if constexpr (!Impl::spmv_general_tpl_avail< + ExecutionSpace, HandleImpl, AMatrix_Internal, + XVector_Internal, YVector_Internal>() && + Impl::spmv_general_tpl_avail< + ExecutionSpace, HandleImpl, AMatrix_Internal, + XVector_SubInternal, YVector_SubInternal>()) { + if (x.extent(1) == size_t(1) && x.span_is_contiguous() && + y.span_is_contiguous()) { + XVector_SubInternal xsub(x.data(), x.extent(0)); + YVector_SubInternal ysub(y.data(), y.extent(0)); + spmv(space, handle->get_impl(), mode, alpha, A, xsub, beta, ysub); + return; + } + } + } + XVector_Internal x_i(x); YVector_Internal y_i(y); diff --git a/sparse/src/KokkosSparse_spmv_handle.hpp b/sparse/src/KokkosSparse_spmv_handle.hpp index 9e7295c72c..a2eecfd1ce 100644 --- a/sparse/src/KokkosSparse_spmv_handle.hpp +++ b/sparse/src/KokkosSparse_spmv_handle.hpp @@ -237,9 +237,8 @@ struct SPMVHandleImpl { ~SPMVHandleImpl() { if (tpl) delete tpl; } - void set_exec_space(const ExecutionSpace& exec) { - if (tpl) tpl->set_exec_space(exec); - } + + ImplType* get_impl() { return this; } /// Get the SPMVAlgorithm used by this handle SPMVAlgorithm get_algorithm() const { return this->algo; } diff --git a/sparse/tpls/KokkosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp b/sparse/tpls/KokkosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp index 9c844ff910..83af85e2ee 100644 --- a/sparse/tpls/KokkosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp +++ b/sparse/tpls/KokkosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp @@ -48,6 +48,7 @@ inline void spmv_bsr_mkl(Handle* handle, sparse_operation_t op, Scalar alpha, if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for MKL BSR"); + subhandle->set_exec_space(exec); } else { // Use the default execution space instance, as classic MKL does not use // a specific instance. @@ -127,6 +128,7 @@ inline void spmv_mv_bsr_mkl(Handle* handle, sparse_operation_t op, Scalar alpha, if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for MKL BSR"); + subhandle->set_exec_space(exec); } else { // Use the default execution space instance, as classic MKL does not use // a specific instance. @@ -392,6 +394,7 @@ void spmv_bsr_cusparse(const Kokkos::Cuda& exec, Handle* handle, if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for cusparse"); + subhandle->set_exec_space(exec); } else { /* create and set the subhandle and matrix descriptor */ subhandle = new KokkosSparse::Impl::CuSparse9_SpMV_Data(exec); @@ -519,6 +522,7 @@ void spmv_mv_bsr_cusparse(const Kokkos::Cuda& exec, Handle* handle, if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for cusparse"); + subhandle->set_exec_space(exec); } else { /* create and set the subhandle and matrix descriptor */ subhandle = new KokkosSparse::Impl::CuSparse9_SpMV_Data(exec); @@ -886,6 +890,7 @@ void spmv_bsr_rocsparse(const Kokkos::HIP& exec, Handle* handle, if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for rocsparse BSR"); + subhandle->set_exec_space(exec); } else { subhandle = new KokkosSparse::Impl::RocSparse_BSR_SpMV_Data(exec); handle->tpl = subhandle; diff --git a/sparse/tpls/KokkosSparse_spmv_mv_tpl_spec_decl.hpp b/sparse/tpls/KokkosSparse_spmv_mv_tpl_spec_decl.hpp index 47b7d47f8e..2ae6bf44f2 100644 --- a/sparse/tpls/KokkosSparse_spmv_mv_tpl_spec_decl.hpp +++ b/sparse/tpls/KokkosSparse_spmv_mv_tpl_spec_decl.hpp @@ -192,6 +192,7 @@ void spmv_mv_cusparse(const Kokkos::Cuda &exec, Handle *handle, if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for cusparse"); + subhandle->set_exec_space(exec); } else { subhandle = new KokkosSparse::Impl::CuSparse10_SpMV_Data(exec); handle->tpl = subhandle; diff --git a/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp b/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp index 926d201a52..9303b0d96a 100644 --- a/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp +++ b/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp @@ -111,6 +111,7 @@ void spmv_cusparse(const Kokkos::Cuda& exec, Handle* handle, const char mode[], if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for cusparse"); + subhandle->set_exec_space(exec); } else { subhandle = new KokkosSparse::Impl::CuSparse10_SpMV_Data(exec); handle->tpl = subhandle; @@ -155,6 +156,7 @@ void spmv_cusparse(const Kokkos::Cuda& exec, Handle* handle, const char mode[], if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for cusparse"); + subhandle->set_exec_space(exec); } else { /* create and set the subhandle and matrix descriptor */ subhandle = new KokkosSparse::Impl::CuSparse9_SpMV_Data(exec); @@ -423,6 +425,7 @@ void spmv_rocsparse(const Kokkos::HIP& exec, Handle* handle, const char mode[], if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for rocsparse CRS"); + subhandle->set_exec_space(exec); } else { subhandle = new KokkosSparse::Impl::RocSparse_CRS_SpMV_Data(exec); handle->tpl = subhandle; @@ -591,6 +594,8 @@ inline void spmv_mkl(Handle* handle, sparse_operation_t op, Scalar alpha, MKLScalar* y_mkl = reinterpret_cast(y); if (handle->is_set_up) { subhandle = dynamic_cast(handle->tpl); + // note: classic mkl only runs on synchronous host exec spaces, so no need + // to call set_exec_space on the subhandle here if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for MKL CRS"); @@ -757,6 +762,7 @@ inline void spmv_onemkl(const execution_space& exec, Handle* handle, if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for OneMKL CRS"); + subhandle->set_exec_space(exec); } else { subhandle = new OneMKL_SpMV_Data(exec); handle->tpl = subhandle; From 9a1856dc6233f52e4cb5897beef8d22e3b1ea89b Mon Sep 17 00:00:00 2001 From: brian-kelley Date: Thu, 4 Apr 2024 11:45:37 -0600 Subject: [PATCH 10/19] Fix #2167: classic MKL doesn't use space instance (#2168) --- sparse/tpls/KokkosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp | 6 ++++-- sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp | 4 ++-- 2 files changed, 6 insertions(+), 4 deletions(-) diff --git a/sparse/tpls/KokkosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp b/sparse/tpls/KokkosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp index 83af85e2ee..e867038842 100644 --- a/sparse/tpls/KokkosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp +++ b/sparse/tpls/KokkosSparse_spmv_bsrmatrix_tpl_spec_decl.hpp @@ -48,7 +48,8 @@ inline void spmv_bsr_mkl(Handle* handle, sparse_operation_t op, Scalar alpha, if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for MKL BSR"); - subhandle->set_exec_space(exec); + // note: classic mkl only runs on synchronous host exec spaces, so no need + // to call set_exec_space on the subhandle here } else { // Use the default execution space instance, as classic MKL does not use // a specific instance. @@ -128,7 +129,8 @@ inline void spmv_mv_bsr_mkl(Handle* handle, sparse_operation_t op, Scalar alpha, if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for MKL BSR"); - subhandle->set_exec_space(exec); + // note: classic mkl only runs on synchronous host exec spaces, so no need + // to call set_exec_space on the subhandle here } else { // Use the default execution space instance, as classic MKL does not use // a specific instance. diff --git a/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp b/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp index 9303b0d96a..e3f88e6e11 100644 --- a/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp +++ b/sparse/tpls/KokkosSparse_spmv_tpl_spec_decl.hpp @@ -594,11 +594,11 @@ inline void spmv_mkl(Handle* handle, sparse_operation_t op, Scalar alpha, MKLScalar* y_mkl = reinterpret_cast(y); if (handle->is_set_up) { subhandle = dynamic_cast(handle->tpl); - // note: classic mkl only runs on synchronous host exec spaces, so no need - // to call set_exec_space on the subhandle here if (!subhandle) throw std::runtime_error( "KokkosSparse::spmv: subhandle is not set up for MKL CRS"); + // note: classic mkl only runs on synchronous host exec spaces, so no need + // to call set_exec_space on the subhandle here } else { // Use the default execution space instance, as classic MKL does not use // a specific instance. From 95310c8e15823377389f6c5705e094d296ab4f8c Mon Sep 17 00:00:00 2001 From: Nathan Ellingwood Date: Wed, 24 Apr 2024 19:35:47 -0600 Subject: [PATCH 11/19] Merge pull request #2186 from ndellingwood/osx-ci-master workflows/osx.yml: test against most recent kokkos tag (cherry picked from commit c099ceae64c0095fbd98df4c9c944aed80ba49c3) --- .github/workflows/osx.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/osx.yml b/.github/workflows/osx.yml index 944807b032..9f05579fa5 100644 --- a/.github/workflows/osx.yml +++ b/.github/workflows/osx.yml @@ -58,7 +58,7 @@ jobs: uses: actions/checkout@v4 with: repository: kokkos/kokkos - ref: ${{ github.base_ref }} + ref: 4.3.00 path: kokkos - name: configure_kokkos From 5fc7e229a763d13e44305eb8ae1defd0c3951fd5 Mon Sep 17 00:00:00 2001 From: Nathan Ellingwood Date: Wed, 1 May 2024 17:35:36 -0600 Subject: [PATCH 12/19] Update changelog for 4.3.01 --- CHANGELOG.md | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 4e6da70740..63d305dd04 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,5 +1,15 @@ # Change Log +## [4.3.01](https://github.com/kokkos/kokkos-kernels/tree/4.3.01) +[Full Changelog](https://github.com/kokkos/kokkos-kernels/compare/4.3.00...4.3.01) + +### Bug Fixes: +- sparse: block spiluk fixes [\#2172](https://github.com/kokkos/kokkos-kernels/pull/2172) +- magma: tpl interaction fixes [\#2176](https://github.com/kokkos/kokkos-kernels/pull/2176), [\#2178](https://github.com/kokkos/kokkos-kernels/pull/2178), [\#2181](https://github.com/kokkos/kokkos-kernels/pull/2181) +- trsv: Add early return if numRows == 0 in trsv to avoid integer divide-by-zero error [\#2180](https://github.com/kokkos/kokkos-kernels/pull/2180) +- blas tpl: resolve potential duplicate symbol [\#2183](https://github.com/kokkos/kokkos-kernels/pull/2183) +- spmv: permformance fix, add special path for rank-2 x/y [\#2164](https://github.com/kokkos/kokkos-kernels/pull/2164), [\#2168](https://github.com/kokkos/kokkos-kernels/pull/2168) + ## [4.3.00](https://github.com/kokkos/kokkos-kernels/tree/4.3.00) (2024-03-19) [Full Changelog](https://github.com/kokkos/kokkos-kernels/compare/4.2.01...4.3.00) From 852c5cb6942c72de93d870e8784d00a5dc67a3b1 Mon Sep 17 00:00:00 2001 From: James Foucar Date: Fri, 3 May 2024 14:05:32 -0600 Subject: [PATCH 13/19] BsrMatrix: Fix HostMirror typedef (#2196) It needed to have size_type. --- sparse/src/KokkosSparse_BsrMatrix.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sparse/src/KokkosSparse_BsrMatrix.hpp b/sparse/src/KokkosSparse_BsrMatrix.hpp index db9ef71753..0ce25074b6 100644 --- a/sparse/src/KokkosSparse_BsrMatrix.hpp +++ b/sparse/src/KokkosSparse_BsrMatrix.hpp @@ -362,7 +362,8 @@ class BsrMatrix { typedef SizeType size_type; //! Type of a host-memory mirror of the sparse matrix. - typedef BsrMatrix + typedef BsrMatrix HostMirror; //! Type of the graph structure of the sparse matrix. typedef Kokkos::StaticCrsGraph Date: Fri, 3 May 2024 15:07:51 -0600 Subject: [PATCH 14/19] Update changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 63d305dd04..c848e7cd14 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -9,6 +9,7 @@ - trsv: Add early return if numRows == 0 in trsv to avoid integer divide-by-zero error [\#2180](https://github.com/kokkos/kokkos-kernels/pull/2180) - blas tpl: resolve potential duplicate symbol [\#2183](https://github.com/kokkos/kokkos-kernels/pull/2183) - spmv: permformance fix, add special path for rank-2 x/y [\#2164](https://github.com/kokkos/kokkos-kernels/pull/2164), [\#2168](https://github.com/kokkos/kokkos-kernels/pull/2168) +- BsrMatrix: Fix HostMirror typedef [\#2196](https://github.com/kokkos/kokkos-kernels/pull/2196) ## [4.3.00](https://github.com/kokkos/kokkos-kernels/tree/4.3.00) (2024-03-19) [Full Changelog](https://github.com/kokkos/kokkos-kernels/compare/4.2.01...4.3.00) From 9402f8a4b7c16120653c87c509727b7fd653fd47 Mon Sep 17 00:00:00 2001 From: Carl Pearson Date: Thu, 2 May 2024 08:23:57 -0600 Subject: [PATCH 15/19] Fix macOS docs build (#2190) * Fix docs build * try docs fix * make sphinx available at config time --- .github/workflows/docs.yml | 16 +++++++++++----- 1 file changed, 11 insertions(+), 5 deletions(-) diff --git a/.github/workflows/docs.yml b/.github/workflows/docs.yml index 04a1ba74b2..1646300e81 100644 --- a/.github/workflows/docs.yml +++ b/.github/workflows/docs.yml @@ -16,9 +16,11 @@ jobs: - name: Install Dependencies run: | brew install doxygen - python3 -m pip install sphinx -v "sphinx==6.2.1" - python3 -m pip install breathe - python3 -m pip install sphinx-rtd-theme + python3 -m venv .venv + . .venv/bin/activate + pip install sphinx -v "sphinx==6.2.1" + pip install breathe + pip install sphinx-rtd-theme sphinx-build --version doxygen --version @@ -52,8 +54,10 @@ jobs: working-directory: kokkos/build run: make -j2 install + # sphinx needs to be available at configure time for the target to be generated - name: configure_kokkos_kernels run: | + . .venv/bin/activate mkdir -p kokkos-kernels/{build,install} cd kokkos-kernels/build cmake \ @@ -81,5 +85,7 @@ jobs: fi - name: build_kokkos_kernels_sphinx - working-directory: kokkos-kernels/build - run: make Sphinx + run: | + . .venv/bin/activate + cd kokkos-kernels/build + make Sphinx From 01cbb989d97fbb37ef49449a341ab9940d397b2c Mon Sep 17 00:00:00 2001 From: Nathan Ellingwood Date: Fri, 3 May 2024 15:08:59 -0600 Subject: [PATCH 16/19] Update changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index c848e7cd14..4f845771c7 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -10,6 +10,7 @@ - blas tpl: resolve potential duplicate symbol [\#2183](https://github.com/kokkos/kokkos-kernels/pull/2183) - spmv: permformance fix, add special path for rank-2 x/y [\#2164](https://github.com/kokkos/kokkos-kernels/pull/2164), [\#2168](https://github.com/kokkos/kokkos-kernels/pull/2168) - BsrMatrix: Fix HostMirror typedef [\#2196](https://github.com/kokkos/kokkos-kernels/pull/2196) +- GA: Fix macOS docs build [\#2190](https://github.com/kokkos/kokkos-kernels/pull/2190) ## [4.3.00](https://github.com/kokkos/kokkos-kernels/tree/4.3.00) (2024-03-19) [Full Changelog](https://github.com/kokkos/kokkos-kernels/compare/4.2.01...4.3.00) From ca09900d98508ddb7e775e0447feafd23a9f9751 Mon Sep 17 00:00:00 2001 From: brian-kelley Date: Tue, 7 May 2024 11:52:55 -0600 Subject: [PATCH 17/19] Update CHANGELOG.md --- CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 4f845771c7..9cb40b5e74 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -8,7 +8,7 @@ - magma: tpl interaction fixes [\#2176](https://github.com/kokkos/kokkos-kernels/pull/2176), [\#2178](https://github.com/kokkos/kokkos-kernels/pull/2178), [\#2181](https://github.com/kokkos/kokkos-kernels/pull/2181) - trsv: Add early return if numRows == 0 in trsv to avoid integer divide-by-zero error [\#2180](https://github.com/kokkos/kokkos-kernels/pull/2180) - blas tpl: resolve potential duplicate symbol [\#2183](https://github.com/kokkos/kokkos-kernels/pull/2183) -- spmv: permformance fix, add special path for rank-2 x/y [\#2164](https://github.com/kokkos/kokkos-kernels/pull/2164), [\#2168](https://github.com/kokkos/kokkos-kernels/pull/2168) +- spmv: permformance fix, add back special path for rank-2 x/y with 1 column [\#2164](https://github.com/kokkos/kokkos-kernels/pull/2164), [\#2168](https://github.com/kokkos/kokkos-kernels/pull/2168) - BsrMatrix: Fix HostMirror typedef [\#2196](https://github.com/kokkos/kokkos-kernels/pull/2196) - GA: Fix macOS docs build [\#2190](https://github.com/kokkos/kokkos-kernels/pull/2190) From 58785c1b02ae31ea5ae7c7c9f8905d8c543876e1 Mon Sep 17 00:00:00 2001 From: Nathan Ellingwood Date: Tue, 7 May 2024 12:33:37 -0600 Subject: [PATCH 18/19] Update version check for 4.3.1 --- CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index e575263e55..45e91a90f0 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -127,13 +127,13 @@ ELSE() IF (NOT KOKKOSKERNELS_HAS_TRILINOS AND NOT KOKKOSKERNELS_HAS_PARENT) # This is a standalone build FIND_PACKAGE(Kokkos REQUIRED) - IF((${Kokkos_VERSION} VERSION_GREATER_EQUAL "4.1.0") AND (${Kokkos_VERSION} VERSION_LESS_EQUAL "4.3.0")) + IF((${Kokkos_VERSION} VERSION_GREATER_EQUAL "4.1.0") AND (${Kokkos_VERSION} VERSION_LESS_EQUAL "4.3.1")) MESSAGE(STATUS "Found Kokkos version ${Kokkos_VERSION} at ${Kokkos_DIR}") IF((${Kokkos_VERSION} VERSION_GREATER "4.3.99")) 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 Kokkos_VERSION 4.1.0, 4.2.0, 4.2.1 or 4.3.0") + MESSAGE(FATAL_ERROR "Kokkos Kernels ${KokkosKernels_VERSION} requires Kokkos_VERSION 4.1.0, 4.2.0, 4.2.1, 4.3.0, or 4.3.1") ENDIF() ENDIF() From 3939ace79cea4835cff729d5363fba9b3b258870 Mon Sep 17 00:00:00 2001 From: Nathan Ellingwood Date: Tue, 7 May 2024 11:16:00 -0600 Subject: [PATCH 19/19] update master_history.txt --- master_history.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/master_history.txt b/master_history.txt index 2207bca133..3e8f8fcbd8 100644 --- a/master_history.txt +++ b/master_history.txt @@ -25,3 +25,4 @@ 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 tag: 4.3.00 date: 04/03/2024 master: afd65f03 release: ebbf4b78 +tag: 4.3.01 date: 05/07/2024 master: 1b0a15f5 release: 58785c1b