Triangular solvers #1193
Open
Triangular solvers #1193
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
adds these triangular solvers: - level based solver - warp inverting solver 'winv' - load-balancing solver 'wvar' - thinned out syncfree solver - block solver
|
label! |
ginkgo-bot
added
mod:all
upsj
reviewed
Feb 8, 2023
| @@ -73,7 +73,7 @@ template <typename ValueType, typename IndexType> | |||
| void generate(std::shared_ptr<const HipExecutor> exec, | |||
| const matrix::Csr<ValueType, IndexType>* matrix, | |||
| std::shared_ptr<solver::SolveStruct>& solve_struct, | |||
| bool unit_diag, const solver::trisolve_algorithm algorithm, | |||
| bool unit_diag, const solver::trisolve_strategy strategy, | |||
There was a problem hiding this comment.
Suggested change
| bool unit_diag, const solver::trisolve_strategy strategy, | |
| bool unit_diag, const solver::trisolve_strategy* strategy, |
| @@ -73,7 +73,7 @@ template <typename ValueType, typename IndexType> | |||
| void generate(std::shared_ptr<const HipExecutor> exec, | |||
| const matrix::Csr<ValueType, IndexType>* matrix, | |||
| std::shared_ptr<solver::SolveStruct>& solve_struct, | |||
| bool unit_diag, const solver::trisolve_algorithm algorithm, | |||
| bool unit_diag, const solver::trisolve_strategy strategy, | |||
There was a problem hiding this comment.
Suggested change
| bool unit_diag, const solver::trisolve_strategy strategy, | |
| bool unit_diag, const solver::trisolve_strategy* strategy, |
| @@ -70,7 +70,7 @@ template <typename ValueType, typename IndexType> | |||
| void generate(std::shared_ptr<const DpcppExecutor> exec, | |||
| const matrix::Csr<ValueType, IndexType>* matrix, | |||
| std::shared_ptr<solver::SolveStruct>& solve_struct, | |||
| bool unit_diag, const solver::trisolve_algorithm algorithm, | |||
| bool unit_diag, const solver::trisolve_strategy strategy, | |||
There was a problem hiding this comment.
Suggested change
| bool unit_diag, const solver::trisolve_strategy strategy, | |
| bool unit_diag, const solver::trisolve_strategy* strategy, |
| @@ -70,7 +70,7 @@ template <typename ValueType, typename IndexType> | |||
| void generate(std::shared_ptr<const DpcppExecutor> exec, | |||
| const matrix::Csr<ValueType, IndexType>* matrix, | |||
| std::shared_ptr<solver::SolveStruct>& solve_struct, | |||
| bool unit_diag, const solver::trisolve_algorithm algorithm, | |||
| bool unit_diag, const solver::trisolve_strategy strategy, | |||
There was a problem hiding this comment.
Suggested change
| bool unit_diag, const solver::trisolve_strategy strategy, | |
| bool unit_diag, const solver::trisolve_strategy* strategy, |
| @@ -32,13 +32,22 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. | |||
|
|
|||
| template <typename ValueType, typename IndexType> | |||
| __device__ __forceinline__ | |||
| std::enable_if_t<!is_complex_s<ValueType>::value, ValueType> | |||
| std::enable_if_t<std::is_floating_point<ValueType>::value, ValueType> | |||
There was a problem hiding this comment.
does this have to do with int not having a definition of is_complex? We will probably run into issues with #1257
Comment on lines
+314
to
+318
| void should_perform_transpose_kernel(std::shared_ptr<const CudaExecutor> exec, | ||
| bool& do_transpose) | ||
| { | ||
| do_transpose = false; | ||
| } |
There was a problem hiding this comment.
this most likely belongs outside the unnamed namespace
Comment on lines
+732
to
+734
| if (levels[row] != sweep) { | ||
| return; | ||
| } |
There was a problem hiding this comment.
does it make sense to sort by levels and instead launch only threads for the rows in that level?
|
|
||
|
|
||
| template <typename IndexType, bool is_upper> | ||
| __global__ void level_generation_kernel(const IndexType* const rowptrs, |
There was a problem hiding this comment.
maybe we can find a way to merge this kernel with the numerical one (same for legacy) with the right generalizations of + and *?
| const auto local_row = rhs; | ||
| const auto row = gid; | ||
|
|
||
| const auto activemask = __activemask(); |
There was a problem hiding this comment.
activemask and Volta don't play nice with each other :-( This may only be a subset of all threads in this branch
|
|
||
| // This is "heavily inspired" by cppreference. | ||
| template <class T> | ||
| __device__ const T* lower_bound(const T* first, const T* const last, |
There was a problem hiding this comment.
we have an implementation in searching.cuh
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This PR adds several triangular solvers. Machinery to either select the solver from a configuration file or to automatically choose a good solver will be included in a future PR. For now, the default remains
sparselib.These solvers are added:
The changes are a little cluttered due to some renaming, e. g.
trisolve_algorithmtotrisolve_strategy.Some minimal refactoring is still required, but in principle the code is ready for review. There are several TODOs still in the code. They are not really pressing, but it should probably be decided what to do about them. Some might just be deleted.