Skip to content

fix(rocSOLVER) Fix race condition in SYEVJ GPU kernel on gfx11 (LLVM Issue #207883) - #12266

Merged
EdDAzevedo merged 12 commits into
developfrom
users/EdDAzevedo/rocm_31219_syevj
Oct 15, 2026
Merged

EdDAzevedo merged 12 commits into
developfrom
users/EdDAzevedo/rocm_31219_syevj

Conversation

@EdDAzevedo

@EdDAzevedo EdDAzevedo commented Sep 17, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

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

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

Pass google tests for all routines that use the Jacobi method on all architectures, including gfx11.

Test Result

All tests passed.

Submission Checklist

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

Restore reference-result validation and add regression coverage for the affected gfx11 sizes.

Get a fresh assessment by requesting another Copilot review.

Pull request overview

Fixes a gfx11 SYEVJ GPU-kernel race by adding synchronization around LDS accesses. The change also removes validation checks from host LAPACK reference helpers.

Changes:

  • Adds synchronization barriers for shared-memory reads and writes.
  • Removes info == 0 assertions from reference routines.
  • Requires regression coverage for affected matrix sizes.
File summaries
File Summary
projects/rocsolver/library/src/lapack/roclapack_syevj_heevj.hpp Adds barriers protecting SYEVJ/HEEVJ LDS accesses.
projects/rocsolver/clients/common/misc/lapack_host_reference.cpp Removes host-reference LAPACK status checks.
Review details

Suppressed comments (4)

projects/rocsolver/clients/common/misc/lapack_host_reference.cpp:8651

  • Removing these assertions from the sytf2/sytrf reference wrappers makes factorization failures (including singular pivots) invisible at the wrapper boundary. The associated tests can then compare results without an explicit reference-status check, weakening failure diagnostics; please preserve the check or return/handle info explicitly.
    ssytf2_(&uploC, &n, A, &lda, ipiv, info);

projects/rocsolver/clients/common/misc/lapack_host_reference.cpp:6331

  • This removes the debug check that the sytrs2 reference solve succeeded. Without replacing it with explicit info handling, a failed CPU reference can be consumed as if it were valid and obscure the real test failure; please preserve the check or propagate the status to the caller.
    ssytrs2_(&uploC, &n, &nrhs, A, &lda, ipiv, B, &ldb, &(work[0]), &info);

projects/rocsolver/clients/common/misc/lapack_host_reference.cpp:8702

  • This removes the debug-time check for sytrf's returned info while the wrapper still writes the factorization outputs. If the CPU factorization fails, the test can proceed with an invalid reference without any diagnostic; preserve validation or propagate and handle the status explicitly.
    ssytrf_(&uploC, &n, A, &lda, ipiv, work, &lwork, info);

projects/rocsolver/library/src/lapack/roclapack_syevj_heevj.hpp:178

  • The check-in SYEVJ/HEEVJ parameterization still samples n=40 and n=67, so it does not exercise the gfx11 small-kernel window around n=48–58 that this synchronization change is intended to fix. Please add regression cases for the affected sizes (or otherwise add a focused test) so this barrier ordering is covered in CI.
    __syncthreads();
  • Files reviewed: 2/2 changed files
  • Comments generated: 1
  • Review effort level: Lite

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

@EdDAzevedo
EdDAzevedo requested a review from a team as a code owner September 17, 2026 20:05
@therock-pr-bot

therock-pr-bot Bot commented Sep 17, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

The added size cases do not reliably reproduce the data-dependent, intermittent race.

Get a fresh assessment by requesting another Copilot review.

Review details
  • Files reviewed: 4/4 changed files
  • Comments generated: 1
  • Review effort level: Balanced

@EdDAzevedo
EdDAzevedo merged commit a56dc4f into develop Oct 15, 2026
161 of 165 checks passed
@EdDAzevedo
EdDAzevedo deleted the users/EdDAzevedo/rocm_31219_syevj branch October 15, 2026 16:38
NguyenNhuDi pushed a commit that referenced this pull request Oct 15, 2026
…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 added a commit that referenced this pull request Oct 15, 2026
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants