Repository navigation
perf(rocsparse): avoid per-row L2 write-back/invalidate on gfx942/gfx950 - #13374
Merged
Merged
Conversation
… factorizations on gfx942/gfx950 On gfx942/gfx950, device-scope acquire/release fences lower to whole-L2 invalidate (buffer_inv sc1) and write-back (buffer_wbl2). The level-scheduled solves and incomplete factorizations issued them once per row. - Solves and factorizations (CSR, CSC, COO, ELL, BSR; sv, sm, ilu0, ic0, ildlt0) on gfx942/gfx950 exchange row results and done flags with device-scope relaxed atomics instead of fences (done_flag_bypass_l2). The result stores complete before the done flag is set. - The analysis of the CSR/CSC/COO/BSR solves and factorizations reduces per workgroup in shared memory and issues a single device-scope atomic, instead of one atomic per row on the same address. Other architectures keep the previous code paths. Co-authored-by: Cursor <cursoragent@cursor.com>
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
- Rename done_flag_bypass_l2 to has_noncoherent_xcd_l2, include __gfx9_4_generic__, and document that the waits rely on the gfx9 wait counters. - Move the store completion wait into wait_store_completion, and emit a workgroup acquire fence in acquire_done_rows so the compiler keeps the result loads after the done flag loads. - Add store_done_row_result for the csrsv and ellsv row results. - Use the __scoped_atomic builtins in the csrsv analysis reduction. - Update stale comments, complete the CHANGELOG entry, and format with clang-format 13. Co-authored-by: Cursor <cursoragent@cursor.com>
- wait_store_completion: the s_waitcnt builtin keeps memory operations on both sides of it; the helper is only valid with the gfx9 wait counters. - Describe the device scope atomics as coherent across the XCD L2s instead of resolved beyond the L2. - Explain the relaxed done-flag store in csrsm and bsrsm, the in-wavefront diagonal exchange in bsric0, and fix the remaining csrilu0 and bsrsm comments. Co-authored-by: Cursor <cursoragent@cursor.com>
The bsrsv kernels for block sizes above 32, the bsrsm kernels, the bsrilu0 kernels for block sizes above 8 and the bsric0 kernels for block sizes above 16 read each dependency result once per block entry. Exchanging the results with device-scope atomics turns every one of these reads into an access that bypasses the L1, which costs more than the L2 write-back and invalidation it replaces. Restore the develop code for these kernels. Co-authored-by: Cursor <cursoragent@cursor.com>
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #13374 +/- ##
===========================================
+ Coverage 71.84% 72.13% +0.29%
===========================================
Files 2876 2869 -7
Lines 475010 473864 -1146
Branches 70171 69892 -279
===========================================
+ Hits 341230 341775 +545
+ Misses 110016 108330 -1686
+ Partials 23764 23759 -5
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
…ilu0 2_8 Dependent block rows read only the diagonal and upper blocks of a block row. Use plain accesses for the remaining blocks, which only the row's own lanes read back, so that they are served from the caches again. Co-authored-by: Cursor <cursoragent@cursor.com>
ntrost57
marked this pull request as draft
October 9, 2026 10:11
In the 2_8, 2_8_unrolled and 9_16 kernels, a block row's own blocks are only read back by the wavefront that stored them, so they need no device-scope loads, and the per-block store wait is not needed either: set_done_flag waits for all stores before the done flag is set. Only the blocks of the dependency rows are read with device-scope loads. Co-authored-by: Cursor <cursoragent@cursor.com>
…vice scope The update of block m in the dot product loop is a plain store, but block m can be a diagonal or upper block that dependent rows read. It is published only because the diagonal and upper blocks are stored again with device scope before the done flag is set. Co-authored-by: Cursor <cursoragent@cursor.com>
In csrsm the explanation of the relaxed done-flag store sat above the old label, and in the small-block bsric0 kernels the note on device-scope loads sat under the old loop label. Merge each into one comment. Co-authored-by: Cursor <cursoragent@cursor.com>
The comment named only spin_loop, but several callers wait with their own relaxed device-scope atomic load loop. State the actual requirement instead. Co-authored-by: Cursor <cursoragent@cursor.com>
ntrost57
marked this pull request as ready for review
October 9, 2026 12:43
On gfx942 and gfx950 the small-block bsric0 kernels store each factored block of the current block row with a device scope store, which bypasses the vector L1. Later iterations read these blocks back as v2 in the dot product, from other lanes of the same wavefront. A plain load there can hit the L1 line that the plain load of the unfactored block allocated earlier and return the old value. The wavefront fences that replaced the per-block wait emit no instructions and do not prevent this. Read these blocks back with load_done_row_data and restore the release_done_row_data wait after storing each block. The initial loads of the own blocks and of the own diagonal block stay plain, since no lane has stored to those blocks yet when they are read. Applies to the 2_8, 2_8_unrolled and 9_16 kernels. Co-authored-by: Cursor <cursoragent@cursor.com>
kliegeois
approved these changes
Oct 9, 2026
jsandham
reviewed
Oct 9, 2026
jsandham
reviewed
Oct 9, 2026
jsandham
approved these changes
Oct 9, 2026
csrilu0 is right-looking: processing column j updates later entries of the current row, including the diagonal, with relaxed device scope stores. On gfx942 and gfx950 the next column loaded its entry before the dependency spin loop, and the diagonal was loaded after the loop, both without waiting for those stores to complete. Load the entry of the current column after the spin loop and wait for the stores of the wavefront first, and wait before the diagonal reload. The spin loop's flag load has usually drained the stores already, so the wait is mostly free. Other architectures keep the previous code. csric0 and csrildlt0 are left-looking and only store the current entry. They reload earlier entries of their row only after the spin loop and keep the diagonal contributions in registers, so they need no change. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
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
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.
Motivation
On gfx942 and gfx950, each XCD has its own L2 cache, and the L2 caches are not coherent with each other. Agent-scope acquire and release fences therefore lower to a full L2 invalidate (
buffer_inv sc1) and a full L2 write-back (buffer_wbl2). The level-scheduled triangular solves and incomplete factorizations issue these fences for every row: the producer before it sets the row's done flag, and the consumer after it has waited for the flag. As a result, these routines spend most of their time on cache maintenance on these architectures.The analysis stage has a second issue. Every row updates the maximum row length with its own device-scope atomic on a single address, and on these architectures these atomics serialize across the XCDs.
Technical details
Solve and factorization kernels. On gfx942, gfx950 and gfx9-4-generic (
has_noncoherent_xcd_l2), row results and done flags are exchanged through device-scope relaxed atomic loads and stores, which carrysc1and are coherent across the XCD L2s. They replace plain accesses ordered by agent-scope fences.__builtin_amdgcn_s_waitcnt(0)behind a workgroup release fence) before it stores the done flag. Nobuffer_wbl2.buffer_inv.The new helpers are in
rocsparse_common.hpp:acquire_done_rowsload_done_row_datastore_done_row_datastore_done_row_result(falls back tonontemporal_storeon other architectures)release_done_row_dataset_done_flagwait_store_completion(gfx9 only; used byrelease_done_row_dataandset_done_flagwhenhas_noncoherent_xcd_l2is true)On all other architectures the helpers compile to the previous code: plain loads and stores, nontemporal stores,
__threadfence(), and the agent-scope acquire fence and release store.Small-block BSR factorizations. In these kernels one wavefront factorizes a block row, so only the blocks that other block rows read need device-scope accesses. The blocks that only the same wavefront reads back use plain loads and stores, because the memory model needs no cache maintenance for coherence between the lanes of one wavefront.
set_done_flagwaits for all stores before it sets the flag.Affected kernels:
bsrsm, bsrsv for block sizes above 32, and the larger-block bsrilu0 and bsric0 kernels keep develop's fence-based synchronization. These kernels reread each dependency's result many times, and every device-scope read bypasses the local L2, so the fence pair per dependency is cheaper there.
Analysis (gfx942, gfx950 and gfx9-4-generic only). Each row updates a per-workgroup maximum row length with a shared-memory atomic. The last row of the workgroup to finish publishes it with a single device-scope atomic. This applies to the analysis of the CSR, CSC, COO, and BSR solves and factorizations, including the analysis stage of
rocsparse_spilu0,rocsparse_spic0, androcsparse_spildlt0. Other architectures keep one device-scope atomic per row.JIRA ID : ROCM-32435