Skip to content

fix(tensilelite): disable TDMPlusLdsBuf (triple LDS buffer) pending race fix - #11654

Merged
mengzcai merged 1 commit into
ROCm:developfrom
mengzcai:users/mengzcai/mxfp8_validator
Sep 7, 2026
Merged

mengzcai merged 1 commit into
ROCm:developfrom
mengzcai:users/mengzcai/mxfp8_validator

Conversation

@mengzcai

@mengzcai mengzcai commented Sep 3, 2026 •

Copy link
Copy Markdown
Contributor

JIRA ID : AIHPBLAS-4459

Motivation

In the ticket : AIHPBLAS-4459. They saw the kernel failed hipblaslt-bench verification run.

Root cause: Triple LDS buffer has an unresolved cross-wave read-after-write race on the rotating LDS

This is a temporary safety gate (marked TEMP in code); the proper fix will restore
triple buffering with the race resolved, then revert this reject.

Technical Details

A fast wave overwrites a buffer a slow wave is still reading. Auto (TDMPlusLdsBuf=-1) falls back to 2 buffers; explicit request (TDMPlusLdsBuf=1) is rejected. Re-enable once the race is fixed.

Test Plan

  • tox (tensilelite): confirm no "0 valid solutions" regression from the reject
  • hipblaslt-test --gtest_filter=matmul: confirm kernels still build & pass
    (triple=-1 falls back to double, no cross-wave race)

Test Result

Testing

Submission Checklist

@therock-pr-bot

therock-pr-bot Bot commented Sep 3, 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/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py; 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 Sep 3, 2026

Copy link
Copy Markdown

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

@codecov

codecov Bot commented Sep 3, 2026 •

Copy link
Copy Markdown

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files
@@           Coverage Diff            @@
##           develop   #11654   +/-   ##
========================================
  Coverage    64.11%   64.11%           
========================================
  Files         2826     2826           
  Lines       469842   469843    +1     
  Branches     69389    69389           
========================================
+ Hits        301195   301197    +2     
  Misses      147536   147536           
+ Partials     21111    21110    -1     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 38.59% <ø> (+<0.01%) ⬆️
TensileLite-Unit 75.55% <100.00%> (+<0.01%) ⬆️
hipBLAS 90.62% <ø> (ø) Carriedforward from fd61638
hipBLASLt 35.22% <ø> (ø)
hipCUB 82.68% <ø> (ø) Carriedforward from fd61638
hipDNN 86.95% <ø> (ø) Carriedforward from fd61638
hipFFT 42.66% <ø> (ø) Carriedforward from fd61638
hipRAND 76.12% <ø> (ø) Carriedforward from fd61638
hipSOLVER 68.92% <ø> (ø) Carriedforward from fd61638
hipSPARSE 86.99% <ø> (ø) Carriedforward from fd61638
rocBLAS 0.28% <ø> (ø) Carriedforward from fd61638
rocFFT 51.25% <ø> (ø) Carriedforward from fd61638
rocRAND 56.91% <ø> (ø) Carriedforward from fd61638
rocSOLVER 76.83% <ø> (ø) Carriedforward from fd61638
rocSPARSE 74.61% <ø> (ø) Carriedforward from fd61638
rocThrust 91.60% <ø> (ø) Carriedforward from fd61638

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines Coverage Δ
...lt/tensilelite/Tensile/SolutionStructs/Solution.py 58.41% <100.00%> (+0.01%) ⬆️

... and 2 files with indirect coverage changes

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
  • 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

…ace fix

Triple LDS buffer has an unresolved cross-wave read-after-write race on the
rotating LDS: a fast wave overwrites a buffer a slow wave is still reading.
Auto (-1) falls back to 2 buffers; explicit request (1) is rejected. Re-enable
once the race is fixed.
@mengzcai
mengzcai force-pushed the users/mengzcai/mxfp8_validator branch from 26f4d3c to 90577b2 Compare September 3, 2026 19:20
@mengzcai
mengzcai merged commit 84aa00a into ROCm:develop Sep 7, 2026
140 of 145 checks passed
@mengzcai
mengzcai deleted the users/mengzcai/mxfp8_validator branch September 7, 2026 04:15
@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.

3 participants