Skip to content

ConjTrans support for blocked serial trsv#2651

Merged
lucbv merged 4 commits into
kokkos:developfrom
yasahi-hpc:refactor-batched-serial-trsv
May 7, 2026
Merged

ConjTrans support for blocked serial trsv#2651
lucbv merged 4 commits into
kokkos:developfrom
yasahi-hpc:refactor-batched-serial-trsv

Conversation

@yasahi-hpc
Copy link
Copy Markdown
Contributor

@yasahi-hpc yasahi-hpc commented May 22, 2025

This PR aims at improving the implementation and testing of serial Trsv. See also #2452

Since blocked version of Trsv relies on Trsm internally, Trsm interface has also been modified.

  • ConjTrans implementation of blocked version
  • Adding a test case for this

Edit on May 5, 2026

  • Improvement in checks. Compile time rank check if the input is a View

Yuuichi Asahi added 3 commits May 5, 2026 13:58
Signed-off-by: Yuuichi Asahi <y.asahi@nr.titech.ac.jp>
Signed-off-by: Yuuichi Asahi <y.asahi@nr.titech.ac.jp>
Signed-off-by: Yuuichi Asahi <y.asahi@nr.titech.ac.jp>
@yasahi-hpc yasahi-hpc force-pushed the refactor-batched-serial-trsv branch from f1e77aa to 49de5c9 Compare May 5, 2026 04:58
Signed-off-by: Yuuichi Asahi <y.asahi@nr.titech.ac.jp>
@yasahi-hpc yasahi-hpc added CI: skip-docs Do not run the documentation checks for this pull request AT2-CI-APPROVAL Approve CI to run at SNL labels May 5, 2026
@lucbv lucbv merged commit e97d153 into kokkos:develop May 7, 2026
26 of 40 checks passed
@ndellingwood
Copy link
Copy Markdown
Contributor

@lucbv looks like something with the changes tripped an issue in ifpack2, from a Serial build with intel oneapi/2023.2.0:

18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/kokkos-kernels/batched/dense/impl/KokkosBatched_Trsv_Serial_Internal.hpp:61:45: error: called object type 'bool' is not a function or function pointer
18:04:56       if (!use_unit_diag) *beta1 = *beta1 / op(A[p * as0 + p * as1]);
18:04:56                                             ^~
18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/packages/ifpack2/src/Ifpack2_BlockTriDiContainer_impl.hpp:4299:7: note: in instantiation of function template specialization 'KokkosBatched::Impl::SerialTrsvInternalLower<KokkosBlas::Algo::Level2::Unblocked>::invoke<bool, double, KokkosBatched::Vector<SIMD<Kokkos::complex<double>>, 4>>' requested here
18:04:56       KOKKOSBATCHED_TRSV_LOWER_NO_TRANSPOSE_INTERNAL_INVOKE(default_mode_type, default_algo_type,
18:04:56       ^
18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/kokkos-kernels/batched/dense/src/KokkosBatched_Trsv_Decl.hpp:146:5: note: expanded from macro 'KOKKOSBATCHED_TRSV_LOWER_NO_TRANSPOSE_INTERNAL_INVOKE'
18:04:56     KOKKOSBATCHED_SERIAL_TRSV_LOWER_NO_TRANSPOSE_INTERNAL_INVOKE(ALGOTYPE, DIAG, M, N, ALPHA, A, AS0, AS1, B, BS);     \
18:04:56     ^
18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/kokkos-kernels/batched/dense/src/KokkosBatched_Trsv_Decl.hpp:84:59: note: expanded from macro 'KOKKOSBATCHED_SERIAL_TRSV_LOWER_NO_TRANSPOSE_INTERNAL_INVOKE'
18:04:56   KokkosBatched::Impl::SerialTrsvInternalLower<ALGOTYPE>::invoke(DIAG::use_unit_diag, false, M, ALPHA, A, AS0, AS1, B, \
18:04:56                                                           ^
18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/packages/ifpack2/src/Ifpack2_BlockTriDiContainer_impl.hpp:4486:7: note: in instantiation of function template specialization 'Ifpack2::BlockTriDiContainerDetails::SolveTridiags<Tpetra::RowMatrix<std::complex<double>>>::solveSingleVector<Kokkos::View<KokkosBatched::Vector<SIMD<Kokkos::complex<double>>, 4> ***, Kokkos::LayoutRight, Kokkos::ScratchMemorySpace<Kokkos::Serial>, Kokkos::MemoryTraits<1>>>' requested here
18:04:56       solveSingleVector(member, blocksize, i0, r0, nrows, v, WW);
18:04:56       ^
18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/kokkos/core/src/Serial/Kokkos_Serial_Parallel_Team.hpp:236:7: note: in instantiation of function template specialization 'Ifpack2::BlockTriDiContainerDetails::SolveTridiags<Tpetra::RowMatrix<std::complex<double>>>::operator()<3>' requested here
18:04:56       m_functor(t, Member(data, ileague, m_league));
18:04:56       ^
18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/kokkos/core/src/Serial/Kokkos_Serial_Parallel_Team.hpp:263:20: note: in instantiation of function template specialization 'Kokkos::Impl::ParallelFor<Ifpack2::BlockTriDiContainerDetails::SolveTridiags<Tpetra::RowMatrix<std::complex<double>>>, Kokkos::TeamPolicy<Kokkos::Serial, Ifpack2::BlockTriDiContainerDetails::SolveTridiags<Tpetra::RowMatrix<std::complex<double>>>::SingleVectorTag<3>>>::exec<Ifpack2::BlockTriDiContainerDetails::SolveTridiags<Tpetra::RowMatrix<std::complex<double>>>::SingleVectorTag<3>>' requested here
18:04:56     this->template exec<typename Policy::work_tag>(
18:04:56                    ^
18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/kokkos/core/src/Kokkos_Parallel.hpp:144:11: note: in instantiation of member function 'Kokkos::Impl::ParallelFor<Ifpack2::BlockTriDiContainerDetails::SolveTridiags<Tpetra::RowMatrix<std::complex<double>>>, Kokkos::TeamPolicy<Kokkos::Serial, Ifpack2::BlockTriDiContainerDetails::SolveTridiags<Tpetra::RowMatrix<std::complex<double>>>::SingleVectorTag<3>>>::execute' requested here
18:04:56   closure.execute();
18:04:56           ^
18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/packages/ifpack2/src/Ifpack2_BlockTriDiContainer_impl.hpp:4901:15: note: in instantiation of function template specialization 'Kokkos::parallel_for<char[45], Kokkos::TeamPolicy<Kokkos::Serial, Ifpack2::BlockTriDiContainerDetails::SolveTridiags<Tpetra::RowMatrix<std::complex<double>>>::SingleVectorTag<3>>, Ifpack2::BlockTriDiContainerDetails::SolveTridiags<Tpetra::RowMatrix<std::complex<double>>>>' requested here
18:04:56       case 3: BLOCKTRIDICONTAINER_DETAILS_SOLVETRIDIAGS(3);
18:04:56               ^
18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/packages/ifpack2/src/Ifpack2_BlockTriDiContainer_impl.hpp:4831:15: note: expanded from macro 'BLOCKTRIDICONTAINER_DETAILS_SOLVETRIDIAGS'
18:04:56       Kokkos::parallel_for("SolveTridiags::TeamPolicy::run<SingleVector>",                                                                            \
18:04:56               ^
18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/packages/ifpack2/src/Ifpack2_BlockTriDiContainer_impl.hpp:5094:22: note: in instantiation of member function 'Ifpack2::BlockTriDiContainerDetails::SolveTridiags<Tpetra::RowMatrix<std::complex<double>>>::run' requested here
18:04:56       solve_tridiags.run(YY, W);
18:04:56                      ^
18:04:56 /home/jenkins/blake-new/workspace/KokkosEco_Trilinos_Blake_OneAPI2023_2_0_ICPX_Serial/Trilinos/packages/ifpack2/src/Ifpack2_BlockTriDiContainer_def.hpp:296:33: note: in instantiation of function template specialization 'Ifpack2::BlockTriDiContainerDetails::applyInverseJacobi<Tpetra::RowMatrix<std::complex<double>>>' requested here
18:04:56     BlockTriDiContainerDetails::applyInverseJacobi<MatrixType>(impl_->A,
...

@yasahi-hpc
Copy link
Copy Markdown
Contributor Author

@ndellingwood
This would fix the issue.
Can you have a look?

Ideally, we do not want users to touch the implementation details, but it is not straightforward to achieve this algorithm without calling these implementation details.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

AT2-CI-APPROVAL Approve CI to run at SNL CI: skip-docs Do not run the documentation checks for this pull request enhancement

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants