Add cuDSS support#1999
Merged
Merged
Conversation
pratikvn
added
mod:cuda
ginkgo-bot
added
reg:build
yhmtsai
previously requested changes
Apr 13, 2026
pratikvn
force-pushed
the
feat/add-cudss
branch
from
cedd1b8 to
f53796c
Compare
pratikvn
force-pushed
the
feat/add-cudss
branch
from
f53796c to
7be0b38
Compare
pratikvn
added
1:ST:ready-for-review
Member
Author
|
cuDSS is now only a part of extensions, but linked to Ginkgo, and also available through the JSON config. |
pratikvn
force-pushed
the
feat/add-cudss
branch
from
25d9de6 to
73b4c43
Compare
yhmtsai
reviewed
Apr 15, 2026
|
|
||
|
|
||
| template <typename ValueType, typename IndexType> | ||
| void CuDss<ValueType, IndexType>::refactorize( |
Member
There was a problem hiding this comment.
Cudss?
CuDss will be cu_dss if we following some trend of naming scheme in Ginkgo
yhmtsai
requested changes
Apr 16, 2026
MarcelKoch
reviewed
Apr 17, 2026
pratikvn
force-pushed
the
feat/add-cudss
branch
from
d9d7e01 to
b7589f6
Compare
yhmtsai
requested changes
Apr 20, 2026
| auto mut_b = const_cast<std::remove_const_t< | ||
| std::remove_pointer_t<decltype(dense_b)>>*>(dense_b); | ||
| for (size_type j = 0; j < nrhs; ++j) { | ||
| mut_b->create_submatrix(span{0, nrows}, span{j, j + 1}) |
Member
There was a problem hiding this comment.
I am surprised at we do not have const version
yhmtsai
requested changes
Apr 27, 2026
Comment on lines
+25
to
+27
| # Absorb into the umbrella ginkgo target (same pattern as ginkgo_cuda etc.) | ||
| # extensions/ is add_subdirectory'd after core/, so the ginkgo target exists. | ||
| target_link_libraries(ginkgo PUBLIC ginkgo_cudss) |
Member
There was a problem hiding this comment.
Should we add this into ginkgo target directly?
It sounds like optional feature such that we include it with Ginkgo.
another approach is only to have ginkgo_cudss target as extension. Users add it when necessary.
Comment on lines
+213
to
+220
| d_wide_output->create_submatrix(gko::span{0, nrows}, gko::span{1, 4}); | ||
|
|
||
| ref_solver->apply(this->input, this->output); | ||
|
|
||
| const auto input_stride_before = strided_input->get_stride(); | ||
| const auto output_stride_before = strided_output->get_stride(); | ||
|
|
||
| cudss_solver->apply(strided_input, strided_output); |
Member
There was a problem hiding this comment.
Suggested change
| d_wide_output->create_submatrix(gko::span{0, nrows}, gko::span{1, 4}); | |
| ref_solver->apply(this->input, this->output); | |
| const auto input_stride_before = strided_input->get_stride(); | |
| const auto output_stride_before = strided_output->get_stride(); | |
| cudss_solver->apply(strided_input, strided_output); | |
| d_wide_output->create_submatrix(gko::span{0, nrows}, gko::span{1, 4}); | |
| const auto input_stride_before = strided_input->get_stride(); | |
| const auto output_stride_before = strided_output->get_stride(); | |
| ref_solver->apply(this->input, this->output); | |
| cudss_solver->apply(strided_input, strided_output); |
| ValueType* x_buf = nullptr; | ||
|
|
||
| if (b_strided) { | ||
| cudaMalloc(&b_buf, nrows * sizeof(ValueType)); |
Member
There was a problem hiding this comment.
cudaMalloc use ginkgo dense allocation
Comment on lines
+294
to
+297
| cudaMemcpy2D( | ||
| b_buf, sizeof(ValueType), dense_b->get_const_values(), | ||
| dense_b->get_stride() * sizeof(ValueType), | ||
| sizeof(ValueType), nrows, cudaMemcpyDeviceToDevice); |
Member
There was a problem hiding this comment.
I would also prefer use ginkgo view copy. At least, this needs to be async version with stream
| } | ||
| if (x_strided) { | ||
| cudaMalloc(&x_buf, nrows * sizeof(ValueType)); | ||
| cudaMemset(x_buf, 0, nrows * sizeof(ValueType)); |
Co-authored-by: Yu-Hsiang Tsai<yhmtsai@gmail.com>
yhmtsai
requested changes
May 4, 2026
Member
yhmtsai
left a comment
yhmtsai
left a comment
There was a problem hiding this comment.
test needs to follow AAA pattern and compare with solution not ref_solver
pratikvn
added
1:ST:ready-to-merge
This file contains hidden or 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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
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.
As our direct solvers are still under development, and we need direct solvers for MICROCARD, I think it would be good to have cuDSS as an option
when usingsolver::Direct.I am open to interface suggestions.
My main constraints were:
Add it toNow extracted intosolver::Directas an additional option, and not as a separate class.extensions, but still available for the user as it is linked as a separate library toginkgo, and also available through the JSON config.refactorize()