From e7b7a94742ad637b3d487aa4f8a547024d9e21a9 Mon Sep 17 00:00:00 2001 From: kliegeois Date: Wed, 21 Oct 2026 21:36:44 +0000 Subject: [PATCH] fix(rocsparse): write the base and use n for csxsldu_preprocess column walks (AISPARSE-906) rocsparse::csxsldu_preprocess_template writes the index base into element 0 of the L and U pointer arrays before the count kernel, which only writes ptr[seq + 1], and the inclusive scan. That write sat in the row-direction case only. In the column direction element 0 kept whatever the caller had left in the arrays, so the scan offset every pointer and the returned lnnz / unnz by that value. The counts were also read back at ptr[m] in both directions. The column walk has n + 1 pointers, so for m < n the counts covered only the first m columns, and for m > n the read went past the last pointer. The base write and the stream synchronization now happen before the direction switch, so both directions start the scan from the base, and the counts are read at ptr[m] for the row walk and ptr[n] for the column walk. The status of the count kernel dispatch, which was dropped, is now returned, so two non-unit diagonals give invalid_value instead of a scan of stale counts. Test: unit_test_host_internal_branches.cpp extends HostInternalBranches.csxsldu_directions, which runs the buffer-size direction combinations and the preprocess walks (row direction with row and column factors, column direction with column factors and a unit or non-unit lower diagonal) on a 2x2 diagonal matrix; the two column walks were missing from develop's version. It adds HostInternalBranches.csxsldu_preprocess_reference, which compares the pointers and counts with a host reference on 2x3, 3x2 and empty 3x0 matrices with strictly lower and upper entries, one-based input, distinct L and U bases, and factors in the other direction (pointers in the scratch buffer), and checks that two non-unit diagonals are rejected. The output arrays start at a -7 sentinel. --- .../unit_test_host_internal_branches.cpp | 203 +++++++++++++++++- .../rocsparse_csxsldu_preprocess.cpp | 28 +-- 2 files changed, 216 insertions(+), 15 deletions(-) diff --git a/projects/rocsparse/clients/unittests/unit_test_host_internal_branches.cpp b/projects/rocsparse/clients/unittests/unit_test_host_internal_branches.cpp index 2d93b9328920..b3fb9fd384e7 100644 --- a/projects/rocsparse/clients/unittests/unit_test_host_internal_branches.cpp +++ b/projects/rocsparse/clients/unittests/unit_test_host_internal_branches.cpp @@ -930,8 +930,8 @@ TEST_F(HostInternalBranches, argdescr_getters_and_verbose) EXPECT_EQ(rocsparse_create_handle(nullptr), rocsparse_status_invalid_pointer); } -// Buffer-size has three direction combinations. Preprocess's two mixed -// diag-type launches are not on the itilu0 path. +// Buffer-size has three direction combinations. Preprocess's column walk and +// the two mixed diag-type launches are not on the itilu0 path. TEST_F(HostInternalBranches, csxsldu_directions) { constexpr auto row = rocsparse_direction_row; @@ -997,6 +997,205 @@ TEST_F(HostInternalBranches, csxsldu_directions) EXPECT_EQ(unnz, 2); EXPECT_EQ(lptr, (std::vector{0, 0, 0})); EXPECT_EQ(uptr, (std::vector{0, 1, 2})); + + // Column direction with column factors: both factors in the input direction, so the + // pointers go to lptr / uptr directly. Every element of lptr / uptr starts at the + // sentinel -7, including element 0. A zero handle buffer makes preprocess allocate + // its scratch. + auto* raw = reinterpret_cast<_rocsparse_handle*>(handle); + size_t saved_buffer = raw->buffer_size; + raw->buffer_size = 0; + lnnz = unnz = -1; + EXPECT_EQ(sldu_preprocess(handle, + column, + rocsparse_diag_type_unit, + column, + rocsparse_diag_type_unit, + column, + scratch.ptr, + &lnnz, + &unnz, + &lptr, + &uptr), + rocsparse_status_success); + raw->buffer_size = saved_buffer; + EXPECT_EQ(lnnz, 0); + EXPECT_EQ(unnz, 0); + EXPECT_EQ(lptr, (std::vector{0, 0, 0})); + EXPECT_EQ(uptr, (std::vector{0, 0, 0})); + + // Same walk with a non-unit lower factor: the diagonal goes to L. + lnnz = unnz = -1; + EXPECT_EQ(sldu_preprocess(handle, + column, + rocsparse_diag_type_non_unit, + column, + rocsparse_diag_type_unit, + column, + scratch.ptr, + &lnnz, + &unnz, + &lptr, + &uptr), + rocsparse_status_success); + EXPECT_EQ(lnnz, 2); + EXPECT_EQ(unnz, 0); + EXPECT_EQ(lptr, (std::vector{0, 1, 2})); + EXPECT_EQ(uptr, (std::vector{0, 0, 0})); + ASSERT_EQ(hipDeviceSynchronize(), hipSuccess); +} + +namespace +{ + struct sldu_case + { + rocsparse_direction dir; + rocsparse_int m, n; + std::vector ptr, ind; + rocsparse_index_base base; + rocsparse_diag_type ldiag; + rocsparse_direction ldir; + rocsparse_index_base lbase; + rocsparse_diag_type udiag; + rocsparse_direction udir; + rocsparse_index_base ubase; + }; + + // Host reference of the L / U pointers that preprocess computes in the input + // direction: one entry per row (row direction) or column (column direction) plus + // one, starting at the factor's index base. At most one diagonal type is non-unit, + // and the diagonal goes to that factor. + void sldu_reference(const sldu_case& c, + std::vector& lptr, + std::vector& uptr) + { + const rocsparse_int nseq = (c.dir == rocsparse_direction_row) ? c.m : c.n; + lptr.assign(1, c.lbase); + uptr.assign(1, c.ubase); + for(rocsparse_int s = 0; s < nseq; ++s) + { + rocsparse_int l = 0, u = 0; + for(rocsparse_int k = c.ptr[s] - c.base; k < c.ptr[s + 1] - c.base; ++k) + { + const rocsparse_int other = c.ind[k] - c.base; + const rocsparse_int row = (c.dir == rocsparse_direction_row) ? s : other; + const rocsparse_int col = (c.dir == rocsparse_direction_row) ? other : s; + if(row > col) + ++l; + else if(row < col) + ++u; + else if(c.ldiag == rocsparse_diag_type_non_unit) + ++l; + else if(c.udiag == rocsparse_diag_type_non_unit) + ++u; + } + lptr.push_back(lptr.back() + l); + uptr.push_back(uptr.back() + u); + } + } +} // namespace + +// Non-square inputs, one-based indices, distinct L / U bases, strictly lower and upper +// entries, and factors in the other direction (their pointers go to the scratch buffer, +// L first, U after it). Every pointer and both counts are compared with the host +// reference; pointer arrays preprocess must not write keep their -7 sentinel. +TEST_F(HostInternalBranches, csxsldu_preprocess_reference) +{ + constexpr auto row = rocsparse_direction_row; + constexpr auto column = rocsparse_direction_column; + constexpr auto unit = rocsparse_diag_type_unit; + constexpr auto nonu = rocsparse_diag_type_non_unit; + constexpr auto zero = rocsparse_index_base_zero; + constexpr auto one = rocsparse_index_base_one; + + // 2x3: (0,0) (0,1) (0,2) / (1,0) (1,1) (1,2), stored by column. + const std::vector p23{0, 2, 4, 6}, i23{0, 1, 0, 1, 0, 1}; + // The same matrix one-based. + const std::vector p23b{1, 3, 5, 7}, i23b{1, 2, 1, 2, 1, 2}; + // 3x2: (0,0) / (1,0) (1,1) / (2,0), stored by column. + const std::vector p32{0, 3, 4}, i32{0, 1, 2, 1}; + // 3x2: (0,0) / (1,0) (1,1) / (2,0) (2,1), stored by row, one-based. + const std::vector r32b{1, 2, 4, 6}, j32b{1, 1, 2, 1, 2}; + + const sldu_case cases[] = { + {column, 2, 3, p23, i23, zero, unit, column, zero, nonu, column, zero}, + {column, 3, 2, p32, i32, zero, nonu, column, zero, unit, column, zero}, + {column, 2, 3, p23b, i23b, one, unit, column, one, nonu, column, zero}, + {column, 2, 3, p23, i23, zero, unit, row, one, nonu, row, one}, + {column, 2, 3, p23, i23, zero, nonu, column, zero, unit, row, one}, + {column, 2, 3, p23b, i23b, one, nonu, row, one, unit, column, zero}, + {row, 3, 2, r32b, j32b, one, unit, row, one, nonu, column, one}, + // 3x0 by column: one pointer and no entries; ind holds a dummy so that the + // device allocation is not empty. + {column, 3, 0, {1}, {1}, one, unit, column, one, nonu, column, zero}, + }; + + int index = 0; + for(const sldu_case& c : cases) + { + SCOPED_TRACE("case " + std::to_string(index++)); + const rocsparse_int nseq = (c.dir == row) ? c.m : c.n; + const rocsparse_int nnz = static_cast(c.ind.size()); + const std::vector sentinel(nseq + 1, -7); + device_vector ptr{c.ptr}; + device_vector ind{c.ind}; + device_vector val{std::vector(nnz, 1.0f)}; + device_vector lptr{sentinel}; + device_vector uptr{sentinel}; + device_vector scratch{std::vector(2 * (nseq + 1), -7)}; + ASSERT_TRUE(ptr.ptr && ind.ptr && val.ptr && lptr.ptr && uptr.ptr && scratch.ptr); + + rocsparse_int lnnz = -1, unnz = -1; + const rocsparse_status status + = rocsparse::csxsldu_preprocess_template( + handle, + c.dir, + c.m, + c.n, + nnz, + ptr.ptr, + ind.ptr, + val.ptr, + c.base, + c.ldiag, + c.ldir, + &lnnz, + lptr.ptr, + c.lbase, + c.udiag, + c.udir, + &unnz, + uptr.ptr, + c.ubase, + scratch.ptr); + EXPECT_EQ(status, rocsparse_status_success); + if(status != rocsparse_status_success) + continue; + + std::vector want_l, want_u; + sldu_reference(c, want_l, want_u); + const bool l_scratch = (c.ldir != c.dir); + const bool u_scratch = (c.udir != c.dir); + const size_t u_offset = l_scratch ? nseq + 1 : 0; + const auto got_l = l_scratch ? to_host(scratch.ptr, nseq + 1) : to_host(lptr); + const auto got_u = u_scratch ? to_host(scratch.ptr + u_offset, nseq + 1) : to_host(uptr); + EXPECT_EQ(got_l, want_l); + EXPECT_EQ(got_u, want_u); + EXPECT_EQ(lnnz, want_l.back() - c.lbase); + EXPECT_EQ(unnz, want_u.back() - c.ubase); + if(l_scratch) + EXPECT_EQ(to_host(lptr), sentinel); + if(u_scratch) + EXPECT_EQ(to_host(uptr), sentinel); + } + + // Both diagonals non-unit is rejected by the count kernel dispatch, in both directions. + for(const rocsparse_direction d : {row, column}) + { + rocsparse_int lnnz = -1, unnz = -1; + EXPECT_EQ(sldu_preprocess(handle, d, nonu, d, nonu, d, nullptr, &lnnz, &unnz), + rocsparse_status_invalid_value); + } ASSERT_EQ(hipDeviceSynchronize(), hipSuccess); } diff --git a/projects/rocsparse/library/src/conversion/rocsparse_csxsldu_preprocess.cpp b/projects/rocsparse/library/src/conversion/rocsparse_csxsldu_preprocess.cpp index 2e1fc0c6fbaf..452a17cd5eaf 100644 --- a/projects/rocsparse/library/src/conversion/rocsparse_csxsldu_preprocess.cpp +++ b/projects/rocsparse/library/src/conversion/rocsparse_csxsldu_preprocess.cpp @@ -159,9 +159,7 @@ namespace rocsparse { case rocsparse_diag_type_non_unit: { - // LCOV_EXCL_START RETURN_IF_ROCSPARSE_ERROR(rocsparse_status_invalid_value); - // LCOV_EXCL_STOP } case rocsparse_diag_type_unit: { @@ -223,19 +221,21 @@ rocsparse_status rocsparse::csxsldu_preprocess_template(rocsparse_handle han : buffer) : uptr_; + RETURN_IF_HIP_ERROR( + rocsparse_hipMemcpyAsync(lptr, &lbase, sizeof(I), hipMemcpyHostToDevice, handle_->stream)); + RETURN_IF_HIP_ERROR( + rocsparse_hipMemcpyAsync(uptr, &ubase, sizeof(I), hipMemcpyHostToDevice, handle_->stream)); + RETURN_IF_HIP_ERROR(rocsparse_hipStreamSynchronize(handle_->stream)); + switch(dir_) { case rocsparse_direction_row: { - RETURN_IF_HIP_ERROR(rocsparse_hipMemcpyAsync( - lptr, &lbase, sizeof(I), hipMemcpyHostToDevice, handle_->stream)); - RETURN_IF_HIP_ERROR(rocsparse_hipMemcpyAsync( - uptr, &ubase, sizeof(I), hipMemcpyHostToDevice, handle_->stream)); - RETURN_IF_HIP_ERROR(rocsparse_hipStreamSynchronize(handle_->stream)); J nblocks = (m_ - 1) / nthreads_per_block + 1; dim3 blocks(nblocks); - rocsparse::csxtril_count_kernel_dispatch( - handle_, blocks, threads, ldiag_, udiag_, m_, ptr_, ind_, base_, lptr, uptr); + RETURN_IF_ROCSPARSE_ERROR( + (rocsparse::csxtril_count_kernel_dispatch( + handle_, blocks, threads, ldiag_, udiag_, m_, ptr_, ind_, base_, lptr, uptr))); RETURN_IF_ROCSPARSE_ERROR(rocsparse::inclusive_scan(handle_, m_, lptr)); RETURN_IF_ROCSPARSE_ERROR(rocsparse::inclusive_scan(handle_, m_, uptr)); break; @@ -245,18 +245,20 @@ rocsparse_status rocsparse::csxsldu_preprocess_template(rocsparse_handle han { J nblocks = (n_ - 1) / nthreads_per_block + 1; dim3 blocks(nblocks); - rocsparse::csxtril_count_kernel_dispatch( - handle_, blocks, threads, udiag_, ldiag_, n_, ptr_, ind_, base_, uptr, lptr); + RETURN_IF_ROCSPARSE_ERROR( + (rocsparse::csxtril_count_kernel_dispatch( + handle_, blocks, threads, udiag_, ldiag_, n_, ptr_, ind_, base_, uptr, lptr))); RETURN_IF_ROCSPARSE_ERROR(rocsparse::inclusive_scan(handle_, n_, lptr)); RETURN_IF_ROCSPARSE_ERROR(rocsparse::inclusive_scan(handle_, n_, uptr)); break; } } + const J nseq = (dir_ == rocsparse_direction_row) ? m_ : n_; RETURN_IF_HIP_ERROR(rocsparse_hipMemcpyAsync( - host_lnnz_, &lptr[m_], sizeof(I), hipMemcpyDeviceToHost, handle_->stream)); + host_lnnz_, &lptr[nseq], sizeof(I), hipMemcpyDeviceToHost, handle_->stream)); RETURN_IF_HIP_ERROR(rocsparse_hipMemcpyAsync( - host_unnz_, &uptr[m_], sizeof(I), hipMemcpyDeviceToHost, handle_->stream)); + host_unnz_, &uptr[nseq], sizeof(I), hipMemcpyDeviceToHost, handle_->stream)); RETURN_IF_HIP_ERROR(rocsparse_hipStreamSynchronize(handle_->stream)); host_lnnz_[0] -= lbase;