Repository navigation
feat(rocprim): tune new device scan params on gfx942 - #13743
Closed
NguyenNhuDi wants to merge 2 commits into
Closed
NguyenNhuDi wants to merge 2 commits into
NguyenNhuDi wants to merge 2 commits into
Conversation
…Issue #207883) (#12266) ## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> This PR fixes a race condition in SYEVJ GPU kernel. The problem showed up on gfx11 architecture. The solution is to add calls to __syncthreads() to avoid race condition hazards in accessing arrays cosine_res[] and sines_diag[] in LDS shared memory. This will address JIRA ROCM-31219 https://amd-hub.atlassian.net/browse/ROCM-31219 ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> The problem was initially identified by github user zpengmei in PR 9135 (now closed) #9135 The issue was initially thought to be related to "[AMDGPU] Miscompiled post-barrier LDS reload in a large-workgroup HIP kernel on gfx11 (silent wrong results) #207883) llvm/llvm-project#207883 Anshil Gandhi (angandhi@amd.com) has identified the source of the problem of race condition where synchronizations are required. This PR adds the calls to __synchthreads() in the GPU kernel for SYEVJ. ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> Pass google tests for all routines that use the Jacobi method on all architectures, including gfx11. ## Test Result <!-- Briefly summarize test outcomes. --> All tests passed. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: Mark Gates <mark.gates@amd.com>
NguyenNhuDi
marked this pull request as ready for review
October 15, 2026 16:45
NguyenNhuDi
changed the title
NguyenNhuDi
force-pushed
the
users/jmahovsk/port-pr12642-rocprim-scan-tuning
branch
from
October 15, 2026 16:48
4532031 to
6141bcf
Compare
✅ 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. |
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
#11702 did add optimizations, but did not include them to the tuning mechanisms. This PR addresses that and tunes said algorithms for MI300X.
This PR should not be merged before #11702. After the aforementioned PR is merged, commit 345fe33 should be dropped from this PR.
JIRA ID : EXSWSTRHPC-330
JIRA ID : EXSWSTRHPC-331
Technical Details
345fe33: #11702 (should be dropped)
96b6d8f: adds the new parameters to the search space
36cee5f: allows tuning with mismatching default parameters
3fd8430: the resulting configs
Test Plan
Tests pass.
Performance improved.
Test Result
CI passes.
Performance verified on MI300X node.
Submission Checklist