fix(rocsparse): remove erroneous __restrict__ qualifiers on shared memory - #10270
Merged
ntrost57 merged 2 commits intoAug 5, 2026
Merged
Conversation
…_blockreduce segmented_blockreduce() in csrmm_device_nnz_split.h takes two pointers into block-shared LDS, and both were marked __restrict__. Each thread reads vals[tid - j], which another thread wrote, so the attribute does not hold. Clang lowers __restrict__ to LLVM noalias. LangRef states that noalias also covers accesses from other threads (llvm/llvm-project#211507). The compiler is therefore free to forward the LDS reads across __syncthreads(). The reduction then drops the contributions of the neighbour threads and returns wrong values. The attribute also buys nothing here. The helper is force-inlined, and its two arguments are distinct __shared__ arrays, so alias analysis already proves that they do not alias. JIRA ID : LCOMPILER-2518 Assisted-by: Claude Opus
✅ 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. |
michaelselehov
changed the title
ronlieb
approved these changes
Aug 1, 2026
ronlieb
left a comment
ronlieb
left a comment
There was a problem hiding this comment.
LGTM, and we need a codeowner to approve
ntrost57
approved these changes
Aug 3, 2026
ntrost57
enabled auto-merge (squash)
August 3, 2026 09:10
kliegeois
approved these changes
Aug 3, 2026
kliegeois
left a comment
kliegeois
left a comment
Contributor
There was a problem hiding this comment.
The changes look good but I am not sure if the changelog modification is in the correct release.
kliegeois
disabled auto-merge
August 4, 2026 13:17
zGoldthorpe
approved these changes
Aug 4, 2026
jsandham
reviewed
Aug 4, 2026
jsandham
approved these changes
Aug 4, 2026
kliegeois
enabled auto-merge (squash)
August 4, 2026 13:55
dso-amd
reviewed
Aug 4, 2026
dso-amd
approved these changes
Aug 4, 2026
dso-amd
left a comment
dso-amd
left a comment
Contributor
There was a problem hiding this comment.
LGTM (see my comments though on the changelog entry).
Co-authored-by: Daniel So <dso@amd.com>
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #10270 +/- ##
===========================================
- Coverage 69.62% 69.62% -0.00%
===========================================
Files 2741 2741
Lines 451540 451737 +197
Branches 66531 66559 +28
===========================================
+ Hits 314366 314493 +127
- Misses 116966 117020 +54
- Partials 20208 20224 +16
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
auto-merge was automatically disabled
August 5, 2026 14:13
Invalid email address
jsandham
added a commit
that referenced
this pull request
Aug 12, 2026
## Motivation Audit and fix incorrect usage of `__restrict__` in rocSPARSE kernels. ## Background Clang lowers `__restrict__` to LLVM `noalias`, which is interpreted to exclude accesses **from other threads**, not just other pointers in the same thread. Combined with the AA change "no synchronization effects for never-escaping locals," this lets the compiler forward/hoist a **plain** load across `__syncthreads()` and drop writes made by sibling threads (the root cause of upstream PR #10270). So `__restrict__` is invalid on: (a) two pointers that provably alias, and (b) any buffer used to exchange data between threads via **non-atomic** reads. ## Fixes applied ### 1. Provable pointer aliasing — ITILU0 `device_calculate` `precond/itilu0/rocsparse_csritilu0_async_inplace.cpp` — `device_calculate` took `T* __restrict__ y_` and `const T* __restrict__ ilu0_`, but both call sites pass `y_ = ilu0_ + k` alongside `ilu0_`, so the two `__restrict__` pointers alias by construction. Removed `__restrict__` from `y_` and `ilu0_`. ### 2. LDS cross-thread reductions (identical to PR #10270) Segmented block-reduction helpers that do plain `data[tid] += data[tid+stride]` across `__syncthreads()` on a `__restrict__` pointer into non-escaping `__shared__` memory: - `extra/csrgemm_device.h` — `csrgemm_group_reduce` - `extra/csrgemm_symbolic_device.h` — `csrgemm_symbolic_group_reduce` - `extra/bsrgemm_device.h` — `bsrgemm_group_reduce` ### 3. Shared-memory hash tables (can hang, not just miscompute) `insert_key` / `insert_pair` / `insert_pair_rxc` do a plain `const I temp = table[hash];` in a loop that must observe other threads' `atomic_cas` writes; `__restrict__` on `table` lets the compiler hoist the load out of the loop. Removed `__restrict__` from `table` (kept it on the atomic-only `data` param) in: - `extra/csrgemm_device.h`, `extra/csrgemm_symbolic_device.h`, `extra/csrgemm_numeric_device.h`, `extra/bsrgemm_device.h` ### 4. Cross-thread producer/consumer on global scratch - **`workspace_B`** (multi-pass GEMM chunk buffer): one lane writes `workspace_B[j]`, all lanes plain-read it next chunk across `__syncthreads()`. Removed `__restrict__` in the 4 device headers and the 4 matching kernel-wrapper `.cpp` signatures (`rocsparse_csrgemm_calc.cpp`, `rocsparse_csrgemm_symbolic_calc.cpp`, `rocsparse_csrgemm_numeric_calc.cpp`, `rocsparse_bsrgemm_calc.cpp`). - **ITILU0 async in-place factors**: `lval_/uval_/dval_` in `csritilu0x_async.cpp` (`kernel_correction`, `kernel_correction_no_norm`) and `ilu0_` in `csritilu0_async_inplace.cpp` (`kernel_calculate`, `kernel_calculate_coo`) are read by some threads while written by others in the same launch (no ping-pong snapshot). Removed `__restrict__` and annotated the kernels.
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
7 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.
Motivation
segmented_blockreduce()incsrmm_device_nnz_split.htakes two pointers into block-shared LDS, and both were marked__restrict__. Each thread readsvals[tid - j], which another thread wrote. The attribute therefore does not hold.Clang lowers
__restrict__to LLVMnoalias. LangRef states thatnoaliasalso covers accesses from other threads, see llvm/llvm-project#211507. The compiler is free to forward the LDS reads across__syncthreads(). The reduction then drops the contributions of the neighbour threads and returns wrong values.The attribute also buys nothing here. The helper is force-inlined, and its two arguments are distinct
__shared__arrays, so alias analysis already proves that they do not alias.JIRA ID : LCOMPILER-2518
Technical Details
This is analogous to #10229 and #9629.
The failure appears with a compiler that contains c3628c7f125b "Reapply [AA] No synchronization effects for never-escaping identified local". That commit tells alias analysis that a synchronizing operation cannot affect an object that never escapes the function. An LDS array declared inside a kernel is such an object. Together with the invalid
__restrict__, this lets the compiler drop the cross-thread partial sums.I confirmed both sides on MI300X (gfx942):
__restrict__present__restrict__removed__restrict__presentTest Plan
Build and run, on MI300X:
The failing cases use the Chebyshev4 matrix,
transA = T, andf64_r.Test Result
Before:
With this patch, all 48 tests pass.
Submission Checklist