fix(ck): [CK] LCOPMILER-2487: Remove erroneous __restrict__ qualifier - #10229
Merged
Conversation
✅ 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. |
ronlieb
approved these changes
Jul 31, 2026
1 task
michaelselehov
approved these changes
Aug 1, 2026
michaelselehov
left a comment
michaelselehov
left a comment
Contributor
There was a problem hiding this comment.
LGTM too. Hopefully one day we can bring the __restrict__ back.
shumway
approved these changes
Aug 3, 2026
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Aug 3, 2026
fix(ck): [CK] LCOPMILER-2487: Remove erroneous `__restrict__` qualifier (#10229) ## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> Memory referenced via `__restrict__`-qualified pointers may not be modified in any way other than through said pointer so long as that pointer is alive. This is ambiguated in the context of multiple threads, but insofar as it is modeled by `noalias` in LLVM, accessing memory through a `__restrict__` pointer in one thread after having it modified by another thread violates this contract. JIRA ID: https://amd-hub.atlassian.net/browse/LCOMPILER-2487 ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> This is analogous to #9629. The semantics for LLVM's `noalias` between threads was clarified in llvm/llvm-project#211507 to also prohibit modifications through the same pointer from other threads. Unless `__restrict__` adopts a weaker guarantee in the future, `p_shared_block` is in violation of this contract. ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> Build and run: `test_gemm_splitk` ## Test Result <!-- Briefly summarize test outcomes. --> Before: ``` [----------] Global test environment tear-down [==========] 32 tests from 8 test suites ran. (77718 ms total) [ PASSED ] 29 tests. [ FAILED ] 3 tests, listed below: [ FAILED ] TestGemmSplitK_MK_NK/0.SmallM, where TypeParam = std::tuple<_Float16,_Float16,_Float16> [ FAILED ] TestGemmSplitK_MK_NK/0.MidLargeM, where TypeParam = std::tuple<_Float16,_Float16,_Float16> [ FAILED ] TestGemmSplitK_MK_NK/0.Regular, where TypeParam = std::tuple<_Float16,_Float16,_Float16> 3 FAILED TESTS ``` With this patch, all tests pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
ntrost57
added a commit
that referenced
this pull request
Aug 5, 2026
…mory (#10270) ## Motivation `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. The attribute therefore does not hold. Clang lowers `__restrict__` to LLVM `noalias`. LangRef states that `noalias` also 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): | compiler | source | result | | --- | --- | --- | | current weekly | `__restrict__` present | 40 pass, 8 fail | | current weekly | `__restrict__` removed | 48 pass | | same commit, AA change disabled | `__restrict__` present | 48 pass | ## Test Plan Build and run, on MI300X: ``` hipsparse-test --gtest_filter='*nightly/spmm_csc.generic*' ``` The failing cases use the Chebyshev4 matrix, `transA = T`, and `f64_r`. ## Test Result Before: ``` [==========] 48 tests from 1 test suite ran. [ PASSED ] 40 tests. [ FAILED ] 8 tests ``` With this patch, all 48 tests pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests --------- Co-authored-by: Nico <31079890+ntrost57@users.noreply.github.com> Co-authored-by: Daniel So <dso@amd.com>
1 task
shumway
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Aug 18, 2026
fix(ck): [CK] LCOPMILER-2487: Remove erroneous `__restrict__` qualifier ## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> Memory referenced via `__restrict__`-qualified pointers may not be modified in any way other than through said pointer so long as that pointer is alive. This is ambiguated in the context of multiple threads, but insofar as it is modeled by `noalias` in LLVM, accessing memory through a `__restrict__` pointer in one thread after having it modified by another thread violates this contract. JIRA ID: https://amd-hub.atlassian.net/browse/LCOMPILER-2487 ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> This is analogous to #9629. The semantics for LLVM's `noalias` between threads was clarified in llvm/llvm-project#211507 to also prohibit modifications through the same pointer from other threads. Unless `__restrict__` adopts a weaker guarantee in the future, `p_shared_block` is in violation of this contract. ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> Build and run: `test_gemm_splitk` ## Test Result <!-- Briefly summarize test outcomes. --> Before: ``` [----------] Global test environment tear-down [==========] 32 tests from 8 test suites ran. (77718 ms total) [ PASSED ] 29 tests. [ FAILED ] 3 tests, listed below: [ FAILED ] TestGemmSplitK_MK_NK/0.SmallM, where TypeParam = std::tuple<_Float16,_Float16,_Float16> [ FAILED ] TestGemmSplitK_MK_NK/0.MidLargeM, where TypeParam = std::tuple<_Float16,_Float16,_Float16> [ FAILED ] TestGemmSplitK_MK_NK/0.Regular, where TypeParam = std::tuple<_Float16,_Float16,_Float16> 3 FAILED TESTS ``` With this patch, all tests pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
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
4 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
Memory referenced via
__restrict__-qualified pointers may not be modified in any way other than through said pointer so long as that pointer is alive.This is ambiguated in the context of multiple threads, but insofar as it is modeled by
noaliasin LLVM, accessing memory through a__restrict__pointer in one thread after having it modified by another thread violates this contract.JIRA ID: https://amd-hub.atlassian.net/browse/LCOMPILER-2487
Technical Details
This is analogous to #9629.
The semantics for LLVM's
noaliasbetween threads was clarified in llvm/llvm-project#211507 to also prohibit modifications through the same pointer from other threads. Unless__restrict__adopts a weaker guarantee in the future,p_shared_blockis in violation of this contract.Test Plan
Build and run:
test_gemm_splitkTest Result
Before:
With this patch, all tests pass.
Submission Checklist