[newling]

Summary

Relands the rocjitsu hipBLASLt race checks from #9363 after the temporary revert in #9938.

The original motivation, design, workload details, limitations, and follow-up work remain documented in #9363. This reland initially covers the existing Linux gfx94X/gfx942 and gfx950 target builds.

The reland adds the following safeguards and scope adjustments:

1. Uses the tensilelite row marker and an explicit changed-path marker so the sidecar runs only when projects/hipblaslt/** selected the final merged product row.
2. Limits the initial race-check coverage to gfx94X/gfx942 and gfx950.
3. Pins rocm-systems to bded72b, ensuring both targets use the same tested rocjitsu revision and exact target-specific KMD configs.
4. Adds continue-on-error: true so rocjitsu failures remain visible without failing the required CI summary during burn-in.
5. Adds the fetched compiler artifact's host-runtime directory to LD_LIBRARY_PATH, allowing tensilelite-client to load libomp.so.
6. Keeps the race sidecar presubmit-only; it is intentionally not enabled in the nightly workflow.

Commit structure

The commits are organized for review:

1. Mechanically reapply ci(hipblaslt): rocjitsu race checks for gfx942/gfx950/gfx1151 #9363.
2. Apply the tensilelite row-selection fix.
3. Add explicit hipBLASLt changed-path selection.
4. Pin rocm-systems and make the check advisory.
5. Add the OpenMP runtime path.
6. Keep the sidecar presubmit-only.
7. Narrow initial coverage to gfx94X/gfx942 and gfx950, simplify the pinned rocjitsu integration, and add explicit provider-row regression coverage.

The first commit is a mechanical restoration; the review-specific changes are isolated in commits 2–7.

JIRA ID: AIHPBLAS-3605

[therock-pr-bot]

✅ 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]

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

[tony-davis]

Thanks for the careful reland. The selection-logic fixes here are solid: the tensilelite token plus the explicit matrix boolean is the right way to stop the old contains('hipblaslt') misfire onto hipblasltprovider, and the new unit tests in therock_matrix_test.py and therock_configure_ci_test.py actually cover that behavior. The rocm-systems pin and the continue-on-error burn-in framing are good instincts too.

My blocking concern is where this runs, not what it does.

Adding gfx1151 to per-PR presubmit is a big compute increase for the whole repo. The change to therock-ci.yml flips the default Linux families to gfx94X, gfx950, gfx1151, and that runs on pull_request for every PR. Because the matrix is project-rows by target-bundles, this adds a full gfx1151 Linux build to every project row on every PR, not just hipBLASLt. In this PR's own run that is 4 new gfx1151 Linux builds (rocrand/hipblaslt row ~26m, rocprim/hipcub/rocthrust ~15m, hipkernelprovider ~10m, hipfft/rocfft ~11m), and 3 of those 4 exist only to feed a hipBLASLt-only sidecar. Given our PR volume this is a large amount of extra compute, and right now it is disclosed as a one-line follow-up rather than called out as a default-behavior change.

The placement is also backwards for a burn-in check. The sidecar is advisory (continue-on-error: true, TODO(AIHPBLAS-3605) to remove after it is reliably green), which is exactly the kind of thing that belongs on nightly first, not on every contributor's PR. As written it is the opposite: it runs on every presubmit, and nightly (therock-ci-nightly.yml) never runs it at all (nightly does not pass run_rocjitsu_race_check, and its Linux families stay gfx94X, gfx950). So we pay for it everywhere except the one place where noise would be cheap, and we get zero nightly coverage of the new check.

Recommendation: target the MultiArch Nightly, not the older single-arch CI. We are moving to MultiArch CI, so new instrumentation like this should land in therock-multi-arch-ci-nightly.yml rather than editing the older therock-ci.yml / therock-ci-linux.yml path. That keeps the experimental check on nightly during burn-in, avoids adding a new runner class to per-PR CI, and puts it on the lane we are actually investing in.

If we do end up touching nightly (or the multi-arch nightly), this PR needs evidence that nightly still runs correctly. Please include a link to a manual workflow run of the affected nightly showing it green, since presubmit CI on this PR does not exercise the nightly path.

To be clear on the smaller stuff, none of that is blocking: the inline TensileLite YAMLs and the GEMM-space expansion already have TODOs, and the second-caller and test_runs_on gating checked out as non-breaking for nightly. The one thing I would want to gate merge on is moving this off per-PR single-arch CI and onto the MultiArch Nightly, with a green nightly run linked here.

[BrianHarrisonAMD]

LGTM with a minor testing nit.

[newling]

Adding gfx1151 to per-PR presubmit is a big compute increase for the whole repo

I can look into either
(1) carving the gfx1151 logic out to be more targeted, and not build and unless tested. This will add code bloat I suspect though, or
(2) dropping gfx1151 in this initial version, and just use gfx950 and gfx942. Long term though, this tool needs to run on multiple architectures.

The placement is also backwards for a burn-in check.

I don't agree with this comment, I think 'burn-in' and 'runs on every PR vs nightly' are orthogonal. Is your concern the overhead (time) of running this test or the possibility that it might fail and block PRs and hold up landing?

In general I have a strong aversion to any testing that does not run on every PR that has the potential to make a test fail, and think nightly testing should only be used when tests are too slow and cannot be made faster.

, and puts it on the lane we are actually investing in.

I don't follow this.

TL;DR I like the idea of MutliArch, because that sounds aligned with running this on different architectures. I don't like the idea of it being nightly as opposed to on every PR. I'll do some investigating though.

[newling]

I am currently working on dropping gfx1151 from this PR altogether (will update PR when ready). That'll eliminate the extra build time being observed. I think this is reasonable: currently there is no gfx1151 testing on PRs for hipblaslt/tensilelite anyway (i.e. no numerical testing or other unit tests) on linux. So forcing through race detection testing for these architectures on linux seems premature. When additional archs are added for 'normal' testing, I can consider adding race detection to accompany.

[tony-davis]

Thanks @newling , that works for me. If we're not adding a full gfx1151 build to every PR, I'm happy with this check running on all hipBLASLt PRs. The job is already gated to hipBLASLt-selected rows (the matrix boolean plus the tensilelite token), and on gfx942/gfx950 it just reuses builds we already do, so the only added cost is the CPU-emulated sidecar itself. That's targeted enough for me.

You're also right that burn-in and per-PR-vs-nightly are orthogonal. My real concern was the per-PR compute from the new gfx1151 build class, not the check running on presubmit. Since it's continue-on-error it can't block landing anyway, so there's no need to touch the nightlies. I'm fine leaving it on every (hipBLASLt) PR.

Dropping gfx1151 for now also seems like the right call given there's no other gfx1151 testing on Linux for hipBLASLt/TensileLite yet. Adding race detection alongside normal testing when more archs come online makes sense. I'll re-review once you push the update, and I'll clear the changes-requested then.

[newling]

@tony-davis thanks for the review above. I've dropped the gfx1151 part of this, so the additional time in CI should be minimal. My understanding is (and I'm going to confirm this) that the new sidecar job runs in parallel to other jobs. As it is fast (the tests run in seconds, their is a build of a relatively small component) it won't add to the total time a PR takes to get through CI (i.e. it won't affect the maximum time over all jobs). If you still have a preference for making this nightly though, happy to discuss.

[newling]

Thanks Tony. Right call to drop gfx1151 for now. I double checked that this new approach (without gfx1151) should add no extra time to a PR getting through CI, an agent confirmed the complete job added in this PR runs in parallel with other normal package tests. And as the complete time for the sidecar is small ( < 4 minutes to run tests, < 1 minute to build the component need ) and does not use the limiting resource (gpu machines) it should not have any effect on total time for a PR.

Commit removing gfx1151: ed4b164

An agent reports that the current run failed for unrelated reasons:

• Nine Windows gfx1151 jobs failed in Driver / GPU sanity check because hipInfo.exe hung before tests.
• One normal Linux hipBLASLt job is stuck in actions/checkout, also before testing.

I'll relaunch the failed tests later when I think this has been resolved

[geomin12]

left a comment. this might require re-organization or a new way to not block rocjitsu race checks

[newling]

As a test that continue-on-error: true works as expected in this PR, I tested using a draft PR #10072, which has an intentional failure to mimic a race. The race checks ran successfully and were then deliberately failed with exit 42; the gfx950 job is visible as failed at https://github.com/ROCm/rocm-libraries/actions/runs/30388307097/job/90386405855. The final summary at https://github.com/ROCm/rocm-libraries/actions/runs/30388307097/job/90424824585
recorded therock-ci-linux: success despite the injected failure; it became red only because therock-ci-windows independently failed. I think this confirms that the nested job-level continue-on-error keeps rocjitsu failures visible without making the required Linux workflow or summary blocking.

I'm going to look into a MI455-style separation still, as it may be useful as a future architectural improvement. I'll also address your other comments @geomin12 and ping when ready for rereview. Thanks for checking this out!

[newling]

@geomin12 could you take another look at this please? No rush, I don't want to land this this week.

[geomin12]

looks right! thanks

[tony-davis]

Thanks for the turnaround. Dropping gfx1151 addresses my blocking concern, so I'm clearing changes-requested.

I checked this is the same code that already ran green: all eight files are byte-identical between the current head 84c46f7 and fd91d7f, so the Aug 3 push was a pure rebase, and at fd91d7f the sidecar passed on both gfx94X and gfx950. I also ran the new unit tests and mutated the source to confirm they bite (removing the tensilelite guard fails test_retrieve_projects_runs_ci_for_workflow_paths). Nice tests.

The current run is red, but it's all "Executing the custom container implementation failed" on the self-hosted runners and develop is green, so I don't think any of it is yours. The side effect is that the hipBLASLt builds died and the race check skipped everywhere, so please re-run before merging so we see the sidecar green at the SHA we land.

Two non-blocking asks. Could you file a ticket for the HPL-init device-fill race that --initialization zero sidesteps? My worry is that "needs independent investigation" in a shell comment quietly becomes permanent. And I'd add default: false to the run_rocjitsu_race_check input so presubmit-only is stated rather than implied.

[estewart08]

This is not always true. In the near future it will be moved to lib/llvm/lib/x86_64-unknown-linux-gnu, but currently it is at lib/llvm/lib. The LLVM_ENABLE_PER_TARGET_RUNTIME_DIR patch in TheRock was reverted due to ASAN build failures.