Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
17 commits
Select commit Hold shift + click to select a range
248e7c0
AISPARSE-677: clamp bsrgemm_calc bucket grids and add grid-stride loops
kliegeois Sep 8, 2026
cf3fdbe
AISPARSE-677: clamp csrgemm_calc/nnz/symbolic/numeric bucket grids an…
kliegeois Sep 8, 2026
cc4a35d
Clamp and grid-stride csrgemm_intermediate_products
kliegeois Sep 10, 2026
4217bc4
fix(rocsparse): widen the bsrgemm/csrgemm calc grid-stride induction …
kliegeois Sep 16, 2026
9b55e9d
style(rocsparse): clang-format the csrgemm calc grid-stride loops (AI…
kliegeois Sep 22, 2026
b5b2ec7
perf(rocsparse): GRID_STRIDE fast path for csrgemm_calc fill kernels …
kliegeois Sep 22, 2026
6570ab5
perf(rocsparse): GRID_STRIDE fast path for bsrgemm/csrgemm numeric+sy…
kliegeois Sep 23, 2026
baf32b4
fix(rocsparse): barrier between grid-stride rows in bsrgemm/csrgemm h…
kliegeois Sep 25, 2026
d26b7d3
Merge branch 'develop' into users/kliegeois/aisparse-677-bsrgemm-calc…
kliegeois Sep 29, 2026
fe8bcea
refactor(rocsparse): use rocsparse_grid.hpp helpers (AISPARSE-677)
kliegeois Sep 29, 2026
8430a11
fix(rocsparse): 64-bit grid-stride counters in csrgemm_intermediate_p…
kliegeois Sep 29, 2026
35a1e8f
docs(rocsparse): name get_grid_size_x as the gemm calc clamp (AISPARS…
kliegeois Sep 29, 2026
bfd4cca
refactor(rocsparse): dispatch GRID_STRIDE through one launch per site…
kliegeois Sep 29, 2026
deb36c4
style(rocsparse): use static_cast instead of C-style casts (AISPARSE-…
kliegeois Sep 29, 2026
76dfa7a
test(rocsparse): forced-clamp unit tests for csrgemm/bsrgemm calc (AI…
kliegeois Sep 29, 2026
b1c1d26
fix(rocsparse): 64-bit induction variable in the block-per-row grid-s…
kliegeois Oct 2, 2026
451f29c
perf(rocsparse): GRID_STRIDE fast path for the csrgemm nnz calc kerne…
kliegeois Oct 2, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions projects/rocsparse/clients/unittests/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -160,6 +160,12 @@ set(ROCSPARSE_UNIT_TEST_DEVICE_SOURCES
unit_test_internal_hostblocks_itilu0.cpp
#10153: distinct anchor to avoid stacked-branch merge conflicts
unit_test_host_extra_bsrgemm.cpp
# --- csrgemm / bsrgemm calc grid.x clamps [AISPARSE-677] -----------------
# Distinct anchor (after the bsrgemm host tests) so this leaf does not conflict
# with the sibling grid-clamp branches that append at the end of the list.
# Drives the public rocsparse_spgemm entry point with
# handle->properties.maxGridSize[0] shrunk; no compile-in TUs.
unit_test_internal_extra_gemm_calc_grids.cpp
# Residual in-scope internal-routine coverage: device __global__ wrappers for
# collective/atomic/hash device helpers, plus host-path primitive helpers.
# unit_test_internal_primitives_host_extras.cpp additionally needs the
Expand Down
72 changes: 72 additions & 0 deletions projects/rocsparse/clients/unittests/unit_test_grid_clamp.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,72 @@
/*! \file */
/* ************************************************************************
* Copyright (C) 2026 Advanced Micro Devices, Inc. All rights Reserved.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in
* all copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
* THE SOFTWARE.
*
* ************************************************************************ */

#pragma once

//
// Test-only helper shared by the large-grid unit tests (AISPARSE-699/700).
//
// The launches hardened by those tickets clamp grid.x with
// rocsparse::get_grid_size_x, to min(maxGridSize[0], (2^32 - 1) / blockDim.x)
// blocks, and grid-stride over whatever the clamp drops. For a 256-thread block
// that clamp is 16,777,215 blocks. The tests shrink maxGridSize[0] for speed
// rather than allocate a problem that large, and run a small problem through
// the same code path. This is the AISPARSE-702 idiom.
//
// A new header rather than an addition to unit_test_utils.hpp: that file is
// shared with the sibling unit-test branches, and a new leaf file cannot
// conflict with them.
//

// Internal handle definition: the complete _rocsparse_handle type is required to
// reach handle->properties.
#include "rocsparse_handle.hpp"

namespace rocsparse_ut
{
// Temporarily shrink the grid.x limit that the launches clamp against, and
// restore it on scope exit so a failed assertion cannot leak the override
// into the next test sharing the fixture's handle.
struct ScopedMaxGridSizeX
{
rocsparse_handle handle;
int saved;

ScopedMaxGridSizeX(rocsparse_handle h, int limit)
: handle(h)
, saved(h->properties.maxGridSize[0])
{
handle->properties.maxGridSize[0] = limit;
}

~ScopedMaxGridSizeX()
{
handle->properties.maxGridSize[0] = saved;
}

ScopedMaxGridSizeX(const ScopedMaxGridSizeX&) = delete;

ScopedMaxGridSizeX& operator=(const ScopedMaxGridSizeX&) = delete;
};
}
Loading
Loading