Skip to content

fix(hipthreads): signal pushCount with hipStreamWriteValue32 to fix gfx12 Windows hang - #13503

Merged
DanielMcIntosh merged 4 commits into
developfrom
users/marsavic/hipthreads-fix-gfx120x-win-hang
Oct 29, 2026
Merged

DanielMcIntosh merged 4 commits into
developfrom
users/marsavic/hipthreads-fix-gfx120x-win-hang

Conversation

@marsavicc

@marsavicc marsavicc commented Oct 12, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

Closes #13487
Closes #13631

On Windows gfx120X, the hipthreads lit tests and the examples hang in TheRock CI. All lit tests that run through hip::wthread time out, and saxpy hits the 30-minute job timeout. The same nightly passes on Windows gfx103X/gfx110X/gfx1151 and on Linux gfx120X.

Technical Details

threading_main is a persistent kernel. It keeps polling cpuWorkQueue.pushCount and exits once the host replaces the -1U value with the final push count. The host wrote that value with hipMemcpyToSymbolAsync. On Windows gfx12, the running kernel never sees it, even after the copy completes. So the kernel never exits, and the next null-stream operation (thrust::copy, hipDeviceSynchronize) waits on it forever.

hipMemcpyToSymbolAsync vs hipStreamWriteValue32:

  • hipMemcpyToSymbolAsync is a data transfer. HIP only guarantees its result is visible to work that runs after it in the stream. A kernel that is already running is not guaranteed to see it.
  • hipStreamWriteValue32 is a stream memory operation meant for signaling work that is already running. It is a single, ordered, aligned 32-bit write.

hipThreads already uses this approach for the work queue: sendToGPU publishes entries with hipStreamWriteValue64 because of a similar Navi4 issue. This PR does the same for pushCount.

Changes in src/hip/thread.cxx:

  • notifyDeviceThereMightNotBeAnyMoreWork and the -1U re-arm in prepDeviceForWork now write pushCount with hipStreamWriteValue32, on the same enqueuing stream as before. The device address is cached.
  • The notify path synchronizes the enqueuing stream so the write is submitted on Windows. This only waits on the enqueuing stream, not on mainStream.
  • The static copy buffer is removed, since the value is now passed directly.

Test Plan

Windows 11, RX 9070 XT (gfx1201), TheRock nightly 10.2.0a20261012:

  • lit (lit -j 1 test/)
  • saxpy step3
  • InOneWeekendRaytracer step4

Test Result

Before After
lit 35 timeouts 89 passed, 0 failed
saxpy hang pass
InOneWeekendRaytracer hang pass

🤖 Generated with Claude Code

@marsavicc
marsavicc requested a review from a team as a code owner October 12, 2026 15:20
@therock-pr-bot

therock-pr-bot Bot commented Oct 12, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ⚠️ Warning Error: Source/code files changed without an accompanying unit test.
Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*).
Current: code file(s) changed: projects/hipthreads/src/hip/thread.cxx; no test file found
🔎 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

therock-pr-bot Bot commented Oct 12, 2026 •

Copy link
Copy Markdown

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

@marsavicc

Copy link
Copy Markdown
Contributor Author

@DanielMcIntosh DanielMcIntosh 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.

Approving this to get the CI green since it seems like it will solve the problem, but I have left some comments that should be addressed in a follow-up PR

Comment thread projects/hipthreads/src/hip/thread.cxx Resolved
Comment thread projects/hipthreads/src/hip/thread.cxx Resolved
Comment thread projects/hipthreads/src/hip/thread.cxx Resolved

@DanielMcIntosh DanielMcIntosh 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.

LGTM

@DanielMcIntosh
DanielMcIntosh enabled auto-merge (squash) October 29, 2026 21:09
@DanielMcIntosh
DanielMcIntosh merged commit c3e9995 into develop Oct 29, 2026
438 of 456 checks passed
@DanielMcIntosh
DanielMcIntosh deleted the users/marsavic/hipthreads-fix-gfx120x-win-hang branch October 29, 2026 21:14
@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

2 participants