[newling]

Motivation

hipBLASLt selects GEMM kernels according to matrix dimensions, numerical types, and operation options. The existing rocjitsu race checks run two small fixed workloads. This PR adds a nightly job that checks kernels reached through a repeatable distribution of GEMM inputs.

The workload samples inputs and lets hipBLASLt select a kernel normally. It complements the approaches that construct several inputs for a selected kernel in #13412 or one input per selected kernel in #13687; it does not force a solution index.

Technical Details

The job runs hipblaslt-bench through CPU emulation with rocjitsu's race detector enabled, so it requires no physical GPU. The detector reports races within one workgroup, such as a register or LDS value read before the required s_waitcnt or s_barrier; it does not see races between workgroups through global memory, such as in stream-K fixups or split-K reductions. See rocjitsu's race detector documentation for details.

A separate sampler draws M, N, and K independently from the multiples of 32 between 32 and 480, with one GEMM per input. It varies six numerical families (FP32, FP16, BF16, FP8, MXFP8, and MXFP4), supported transposes, alpha and beta, bias, and ReLU/GELU output operations. All use FP32 computation. A fixed seed makes the input sequence reproducible across worker counts and builds; normal library selection determines which kernels execute.

Each case must pass the benchmark's CPU-reference numerical checks and show the selected GEMM in the race detector's dispatch log. Races, numerical failures, emulator diagnostics, timeouts, and missing heuristic solutions fail the run. Cases are never replaced. Per-case inputs, commands, selected kernels, logs, and outcomes are retained for diagnosis and replay. The hipBLASLt guide records the sampling policy and local commands.

rocjitsu emulates one GPU target per run, and each target has its own kernels, so the nightly Multi-Arch workflow runs one job per entry in a target matrix. gfx950 (MI350 family) is the only entry for now; adding a target is a matrix entry plus a sampling profile for its number formats. Each job runs 100 sampled GEMMs with seed 0 and eight concurrent cases, using the library and benchmark artifacts from the same nightly run. With a fixed seed, every night checks the same inputs against the latest library, so this is a smoke test rather than a growing search. The existing reusable race-check workflow gains a workload input: smoke (the default) keeps the per-PR check exactly as before, through its unchanged driver and settings, and random-gemms runs this workload through a new driver that builds the same pinned rocjitsu revision, times each stage and records the detector revision. The workflow can also be dispatched manually against a completed run with any count and seed. Each emulated GEMM command is bounded at 900 seconds, so a hang is reported as a timeout while the remaining cases still run. Like the per-PR smoke check, the job is advisory while it settles in: a failure shows on the job and its reports are uploaded, but the nightly run still succeeds, and the job is not part of the required Multi-Arch summary. Nothing notifies on failure, so the guide names an owner who checks the nightly results.

The longer-term plan is a nightly sweep that runs every solution index in the library, with constructed problems instead of random samples. This PR is shaped so that sweep can be added rather than replace it: the sampled-GEMM runner lives in projects/hipblaslt/clients/scripts/rocjitsu, where the sweep runner is planned, shares its option names (--client, --reports, --workers, --library-dir, --target, --continue-on-failure, --rocjitsu-revision) and writes the same kind of summary.md and summary.json report. A new workload then needs a new value of the workload input and a branch in the GEMM driver, while the nightly job, artifact retrieval, rocjitsu build, timings, job summary and Component CI lane stay as they are.

The last hosted run of these 100 inputs, on an earlier revision of this PR, took 66 seconds on an AWS medium runner (5 seconds median and 14 seconds maximum per case). That run also spent about 5 minutes building rocjitsu with limited parallelism; the build now uses all runner cores, which brought it to about 30 seconds. The nightly job should take a few minutes in total.

Each run reports its timings so that changes are easy to follow from day to day. The job summary shows the duration and exit status of each stage (rocjitsu configure and build, the workload), the timeouts in force, and the run totals: cases passed, failed and without a solution, elapsed time, per-case minimum, median and maximum seconds, the five slowest cases, the number of distinct kernels reached, and a table of any failing cases with their reasons. The log prints each case's outcome with its total, conversion and bench seconds. The uploaded reports keep summary.json, cases.md, kernels.json (each kernel and the cases that ran it) and per-case inputs, commands, logs and results.

The runner's unit tests (eight tests, a few seconds) use a stand-in benchmark and emulator. A new Component CI lane, hipBLASLt rocjitsu scripts, runs every test in projects/hipblaslt/clients/scripts/rocjitsu on ubuntu-24.04 when that directory changes, following the existing component-ci-<name>.yml pattern. It rolls up to the required Component CI Summary, so a change that breaks these scripts cannot merge; other PRs skip the lane. The hipBLASLt TESTING.md gate table lists the lane, and its sanitizer table now lists the race detector and its scope.

Test Plan

Run the runner tests for deterministic sampling, concurrent execution, replay, result validation, and failure handling. Run the sampled workload locally under CPU emulation with the pinned rocjitsu revision, and replay deliberately faulty code objects to check failure reporting. Validate the scripts, workflows, and documentation links.

Test Result

The nightly job's driver ran end to end locally with the job's settings, on the gfx950 BLAS artifacts of the November 1 develop nightly fetched by TheRock's artifact script as the job does, with rocjitsu built from the pinned revision: all 100 cases passed with no missing solutions, reaching 56 distinct kernels in 45 seconds with eight workers (3.3 seconds median and 8.4 seconds maximum per case). All eight runner tests pass, including checks of the per-case timings, timeout reporting, distinct-kernel count and summary report, and the new Component CI lane passes on this PR. Earlier local runs of 1,000 inputs also passed every input and reached 149 kernels, and deliberately faulty code objects confirmed that races and numerical errors fail the run. Before per-PR race checks were disabled by #14683 (TheRock-Infra#1160), a hosted run of the 100-input workload on an earlier revision of this PR also passed. The component path check, formatting, workflow lint, and documentation checks pass. The scheduled nightly job will first execute after merge.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Risk level

Low (2/5): adds artifact retrieval, an emulator build, and a short recurring nightly CPU workload to CI. Product code and numerical tolerances are unchanged. The nightly job is advisory and gates nothing. On pull requests, only those that change projects/hipblaslt/clients/scripts/rocjitsu gain a short unit-test lane, which gates Component CI Summary. The per-PR smoke check is unchanged.

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.

[newling]

A few comments

[newling]

Some comments

[newling]

A few more comments

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.

❌ Your project check has failed because the head coverage (70.08%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage.

Additional details and impacted files

@@             Coverage Diff              @@
##           develop   #14123       +/-   ##
============================================
- Coverage    77.71%   70.08%    -7.63%     
============================================
  Files         2904       67     -2837     
  Lines       491016    22641   -468375     
  Branches     74425     3541    -70884     
============================================
- Hits        381556    15866   -365690     
+ Misses       84467     5276    -79191     
+ Partials     24993     1499    -23494     

Flag	Coverage Δ
TensileLite-CPP	?
TensileLite-Unit	?
hipBLAS	?
hipBLASLt	70.08% <ø> (+0.03%)	⬆️
hipCUB	?
hipDNN	?
hipFFT	?
hipRAND	?
hipSOLVER	?
hipSPARSE	?
rocBLAS	?
rocFFT	?
rocRAND	?
rocSOLVER	?
rocSPARSE	?
rocThrust	?
see 2839 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.