[SPARSE] Add support for rocSPARSE backend #544
Conversation
Rbiessy
changed the title
|
Tests log on W6800: amd_w6800_log.txt |
gajanan-choudhary
added
backend
|
Yes, this is in progress! I expect I will need a few days at least. I'm aiming to merge the PR by the end of the year. |
Rbiessy
force-pushed
the
romain/rocsparse
branch
from
af87004
to
c4b952c
Compare
|
@gajanan-choudhary I have updated the PR with recent changes from cuSPARSE. Note that it moved almost all the content from |
| - The same sparse matrix handle cannot be reused for multiple operations | ||
| ``spmm``, ``spmv``, or ``spsv``. Doing so will throw a | ||
| ``oneapi::mkl::unimplemented`` exception. See `#332 | ||
| <https://github.com/ROCm/rocSPARSE/issues/332>`_. |
There was a problem hiding this comment.
Wow, this is quite severe, but seems to be a legitimate issue on rocSPARSE side right now.
There was a problem hiding this comment.
If and when they fix this issue, though, will it be easy for us to make changes (with a version check of course) that correctly performs the operations rather than throwing an unimplemented exception?
There was a problem hiding this comment.
Yes, it's easy to fix on oneMath side. The issue is also referenced in this comment: https://github.com/oneapi-src/oneMKL/pull/544/files#diff-3b8c1c2c71abd54f8f90f43415c2f17b2a7fdb81c2b882c210f3cba56b4679adR63
One would just need to remove the used member, its 2 usages below as well as the mark_used method.
| if (this->format == detail::sparse_format::COO && | ||
| !this->has_matrix_property(matrix_property::sorted)) { |
There was a problem hiding this comment.
In the documentation, docs/domains/sparse_linear_algebra.rst, you've written:
- The CSR format requires the column indices to be sorted within each row.
So shouldn't we be handling both CSR and COO formats here instead of just COO?
| handle_helper.rocsparse_handle_container_mapper_.insert( | ||
| std::make_pair(piPlacedContext_, atomic_ptr)); | ||
|
|
||
| sycl::detail::pi::contextSetExtendedDeleter(*placedContext_, ContextCallback, atomic_ptr); |
There was a problem hiding this comment.
Is it unchanged between the removed PI versus new UR APIs? Isn't there a sycl::detail::ur::contextSetExtendedDeleter somewhere that you need to place under #ifdef ONEAPI_ONEMKL_PI_INTERFACE_REMOVED #else #endif... ?
| auto event = queue.submit([&](sycl::handler& cgh) { | ||
| auto acc = val.template get_access<sycl::access::mode::read_write>(cgh); | ||
| detail::submit_host_task(cgh, queue, [=](sycl::interop_handle ih) { | ||
| if (dvhandle->size != size) { |
There was a problem hiding this comment.
I wonder if we should throw an exception if dvhandle->size < size instead of !=. E.g., users may want to pad thesycl::buffer for some reason with zeros in the end, which wouldn't work with the current if condition.
There was a problem hiding this comment.
Nevermind, this is set_dense_vector_data used for replacing/switching out arrays in an existing handle. The condition is fine then.
There was a problem hiding this comment.
Although maybe it doesn't hurt to add the check both here and in init_xxx_data
| template <typename fpType> | ||
| void init_dense_vector(sycl::queue& queue, dense_vector_handle_t* p_dvhandle, std::int64_t size, | ||
| sycl::buffer<fpType, 1> val) { | ||
| auto event = queue.submit([&](sycl::handler& cgh) { |
There was a problem hiding this comment.
Can we add a check here for sycl::buffer case that checks buffer->size() >= size and throws an exception otherwise?
| CHECK_DESCR_MATCH(spmv_descr, alg, "spmv_optimize"); | ||
|
|
||
| A_handle->mark_used(); | ||
| auto& buffer_size = spmv_descr->temp_buffer_size; |
There was a problem hiding this comment.
Is there a reason why this variable specifically is kept and captured by reference outside compute_functor while other members of spmv_descr are created inside the functor? If it is because buffer_size is used later in an if condition, I'd prefer it if we change this approach and capture everything by [=] in compute_functor, even though it would mean replacing the last buffer_size > 0 with spmv_descr->temp_buffer_size > 0. I know that what you have right now is expected to work in this particular case (because of the condition that spmv_descr must live as long as spmv is being called), it is normally a bad idea to capture variables by reference in SYCL functors that are going to be running asynchronously unless you have an immediate event.wait() (e.g., like what we are doing in spmv_buffer_size() function).
There was a problem hiding this comment.
LGTM other than some minor change requests. This PR was a lot easier to review having reviewed #527. Thanks for the fantastic work!
Description
Add support for the rocSPARSE backend.
Depends on #527 and #532.
Rendered docs: docs.zip
Checklist
All Submissions
Do all unit tests pass locally? amd_log.txt nvidia_log.txt intel_log.txt
Have you formatted the code using clang-format?