Repository navigation
ci(gfx1250): self-build gfx1250-strict per PR - #12983
Merged
tony-davis merged 22 commits intoOct 14, 2026
Merged
Conversation
…e strict artifact; one HW-info log filter for both suites Co-Authored-By: Claude <noreply@anthropic.com>
github-actions
Bot
added
github actions
project: none
⏳ Waiting on Required Checks⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
⏳ Waiting for required checks to start or finish:
|
… mislabelled strict library
Contributor
|
Cancelling the remaining gfx950 jobs on this run so the TheRock submodule bump PRs can be prioritized on the gfx950 queue. This change does not target gfx950, and gfx950 was the only work still outstanding. Jobs that already finished are unchanged. Re-run failed jobs if you need gfx950 results before merge. |
…t built the artifact, harden HW log scrub
1 task done
archana-ramalingam
changed the title
…facts Build hipBLASLt + TensileLite from source against the pinned ROCm SDK nightly, guard that its compiler matches per-PR CI's TheRock pin, and run on pull_request for org-labelled PRs touching hipBLASLt or its shared deps. Per-PR gfx1250-strict artifacts depend on TheRock#8350, and the self-build doesn't wait for the full multi-arch build. The container image lacks the client-build deps, so install libgtest-dev, libgmock-dev (hipblaslt-test) and libzstd-dev (the SDK's LLVM link interface, needed by the TensileLite client). Drop the TensileLite unit tests: they test host code and already run per PR in TheRock CI. Co-Authored-By: Claude <noreply@anthropic.com>
1 of 2 tasks
archana-ramalingam
changed the title
…rkers Both suites ran one test at a time, which left the GPU idle while each test did CPU work. Run hipBLASLt as 4 gtest shards on the same GPU, each scrubbed as it runs, and fail the step if any shard fails. Run TensileLite common with 4 pytest workers; they share the GPU's client lock, so the GPU runs still take turns while kernel compiles overlap them. Co-Authored-By: Claude <noreply@anthropic.com>
Contributor
Author
|
Failing tests are unrelated existing failures/flaky tests. |
Keep the ROCm SDK nightly pinned, and make the compiler check warn instead of fail: when the pinned SDK was built with a different compiler than the one at ci-env's therock-ref, the step summary says so and names the compiler to move the pin to. A failed lookup also only warns. Drop allow_toolchain_mismatch, which only existed to downgrade the old failure. The gate now also runs for PRs that change ci-env/action.yml, so a compiler change shows up on the bump PR, and skips PRs until the repo variable GFX1250_RUNNER_READY is true, so jobs don't queue for an unregistered runner and fail after 24 hours. Fail early unless the job sees exactly one GPU: with one runner per GPU, the gtest shards and pytest workers must stay on the runner's own GPU. Install the SDK and test tools into a venv per job: setup-python's interpreter lives in the self-hosted runner's tool cache, so packages would otherwise carry over to the next PR. Prune pip cache files unread for 14 days, since every nightly adds several GB and pip never evicts. Drop ccache (a cold build takes the same 3 min as a warm one) and the /scratch/blis mount (tasks.py only checks that a BLIS exists, so a pre-staged copy could outlive a version bump; install.sh downloads the pinned one). Run nightly on develop at 03:00 UTC (7 pm PST; GitHub cron is UTC), behind the same runner gate. At the end, hand the workspace and RUNNER_TEMP back to the user that owns _work: the container runs as root, and the runner can't clean root-owned files. Co-Authored-By: Claude <noreply@anthropic.com>
JIRA ID : NA ## Motivation - The gfx1250-strict HW job (#12983) found 12 tests on develop that can't run on gfx1250-strict: 4 TensileLite configs and 8 hipBLASLt MX rows. None of them ran on gfx1250-strict before, because no per-PR CI covers it. - Stacked on #12983 so the HW job runs on this PR; retarget to `develop` once #12983 merges. ## Technical Details - TensileLite: add `skip-gfx1250-strict` to configs that already skip `gfx1250`. A base `skip-gfx1250` mark doesn't fire on gfx1250-strict by design (`test_config_helpers.py`), so these ran there: - `custom_rocblas_gemv_bf16_tn_m{1,2,4}.yaml` (#12329): gfx950 hand-written assembly (`v_dot2c_f32_bf16`, `s_waitcnt`) that doesn't assemble for gfx12. - `streamk/sk_queue_selfclean.yaml` (#10451): SGEMM with a 32x32x2 MFMA tile, rejected as an invalid WMMA configuration (0 solutions). - hipBLASLt `matmul_gtest.yaml` (#13031): `matmul_fast_check_mx_fp8` and `_large_k` run every fp8/bf8 pair on `'(950|1250)'`, which also matches gfx1250-strict. gfx1250 and gfx1250-strict ship transposed-A MX kernels for f8 x f8 (and f8 x f4, f4 x f4) only, so the bf8 pairs fail with `NO solution found!`. The four rows with bf8 inputs get `gpu_arch: '950'` (rows override their entry), so f8 x f8 stays on `'(950|1250)'` and the bf8 pairs run on gfx950 only, like the other MX tests. Entry names and gfx950 test IDs are unchanged. ## Test Plan - 4 non gfx1250 tensilelite tests are skipped on gfx1250 test. - 8 failing hipblaslt tests pass ## Test Result - [ ] gfx1250-strict HW test: - 4 non gfx1250 tensilelite tests are skipped - 8 failing hipblaslt tests pass ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level - Low: 4 skip marks and an arch filter on 8 rows. No library code changes. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com>
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
…ainer Each gfx1250 runner sets ROCR_VISIBLE_DEVICES to its own GPU in its .env. The runner loads .env into its own environment but passes only workflow env into job containers, so the container saw every GPU and the one-GPU check would fail every job. A value-less -e in the container options makes docker copy the variable from the runner. If it is unset, the container still sees every GPU and the one-GPU check fails the job. Co-Authored-By: Claude <noreply@anthropic.com>
The runners are registered now, so PR and nightly jobs no longer queue for a missing runner, which is all the gate was for. The gate also skipped fork PRs from org members, because repository variables are empty on pull_request runs from forks. The on-demand version had no such switch. Gating is back to non-draft, org label and paths for PRs; the nightly and manual dispatch always run. Co-Authored-By: Claude <noreply@anthropic.com>
talumbau
requested changes
Oct 12, 2026
talumbau
left a comment
talumbau
left a comment
Collaborator
There was a problem hiding this comment.
I gave a once through to these Claude generated review comments and they look correct to me. Please take a look and make your own assessment.
4 of 13 tasks
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (49.67%) 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 #12983 +/- ##
========================================
Coverage 71.99% 71.99%
========================================
Files 2879 2879
Lines 478104 478104
Branches 70852 70852
========================================
Hits 344166 344166
Misses 109972 109972
Partials 23966 23966
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
- Build step: set -euo pipefail. apt-get install must succeed; a failed
apt-get update only warns, since install can still use the image's
package lists.
- TensileLite step: run only when setup succeeded (still after a hipBLASLt
build or test failure) and not on cancellation. Install only tox: tox
builds its own env from tox.ini, so the venv installs of
requirements*.txt and -e . were never used. It creates its own results
dir, since the build step may have stopped before making it.
- TensileLite selection: -m "common and not skip-gfx1250-strict". Those
configs were already skipped on this target (config_helpers.py adds a
skip for skip-<arch>), so the same tests run and pass; the expression now
says so. Drop the --deselect of flags/gfx942/test_build_only.py::test_compile:
the test now pins --gpu-targets gfx942 itself.
- Check the gfx1250-strict code objects next to the libhipblaslt that
hipblaslt-test loads, where the runtime looks ({lib dir}/hipblaslt/library),
not always the install tree's copy.
- Mount the pip cache read-only except for the nightly, so a PR job can't
change what later jobs install. Only the nightly prunes it. With a
root-owned cache, pip as root reads it and skips the writes it can't make
(tested).
- HW_SCRUB: rows count columns from the end because the quoted problem
sizes add fields before them. Drop a data row shorter than its header
(e.g. truncated) instead of indexing a negative field, which makes awk
exit 2. Output on the earlier HW logs is unchanged.
- Keep setup-python's LD_LIBRARY_PATH entry when adding the SDK's; without
it every python3 after the SDK step failed to load libpython3.12.so.
- Create the per-job venv with --clear.
Co-Authored-By: Claude <noreply@anthropic.com>
The comment above matmul_fast_check_mx_fp8 said gfx1250 has MX kernels for f8 x f8 only, but it also has f8 x f4 and f4 x f4 MX logic. What it lacks is any bf8 MX kernel. Comment only; no test changes. Co-Authored-By: Claude <noreply@anthropic.com>
The first full CI run failed in two places: - Build: CMake found setup-python's interpreter through the Python_ROOT_DIR/Python3_ROOT_DIR it exports, not the job's venv, so hipblaslt_gentest.py ran without PyYAML. - TensileLite: tox passes PIP_* through, so its pip tried to cache the wheel it builds for tensile in the read-only /pip-cache. Drop actions/setup-python and build the venv from the image's Python 3.12 (apt python3-venv python3-dev), as the local runs on the gfx1250 host do. That also removes the LD_LIBRARY_PATH dependency and the runner tool cache, which persists across jobs. Set PIP_CACHE_DIR=/pip-cache only on the SDK install step. Later pip runs use pip's default cache under the job's HOME. Reword the pip-cache comment: the read-only mount stops a job's build or tests from changing the cache by accident, nothing more. Co-Authored-By: Claude <noreply@anthropic.com>
…atch runs - TensileLite: the first run to reach this step failed in tox's setup. With CC/CXX unset, scikit-build-core sets them to this Python's default compiler (x86_64-linux-gnu-g++), which the image doesn't have. Export the SDK's amdclang/amdclang++ in the step, as the hipBLASLt build does. tox passes CC/CXX through by default, and rocisa's own default is amdclang (cmake/modules/default_amdclang.cmake). - Concurrency: add the event name to the group for non-PR runs, so a manual run on develop and the nightly don't cancel each other. - One-GPU check: require ROCR_VISIBLE_DEVICES (or HIP_VISIBLE_DEVICES) to name exactly one GPU; drop the rocminfo fallback for when neither is set. Every runner sets ROCR_VISIBLE_DEVICES to its own GPU. - Timing: drop wall-clock times inside the test pipelines, before the console and the files, and remove the separate redaction step. It edited the files only after pytest's summary (e.g. "264 passed in 4.92s") had already gone to the public console. - Cleanup: the SDK pin is set once in the job env (the install and the compiler check read it); one apt install in the SDK step; drop PIP_BREAK_SYSTEM_PACKAGES (all pip runs are in a venv); drop shared/rocroller from the checkout and the PR path filter, since the build skips rocRoller; each test step creates its own results directory. The same run passed the build and the hipBLASLt suite: 68,138 tests over 4 shards, 0 failed. Co-Authored-By: Claude <noreply@anthropic.com>
The first full CI run of the TensileLite step still showed run times on the console and in the uploaded log: pytest's "(0:13:09)" (which follows the "in Ns" already removed), tox's "(N seconds)", the Tensile benchmark's "End - Ns" / "_Final Ns", kernel generation's "(N secs elapsed)" and the client CMake's "done (Ns)". Remove them all in the same pipeline, before the scrub and tee. Checked on that run's log: only the "timeout: 600.0s" setting remains. That run otherwise passed: hipBLASLt 68,138 tests over 4 shards, 0 failed; TensileLite 109 passed, 2 xfailed, 1 failed (gsu_gfx1250-strict.yaml, the known GSU multi-buffer race: wrong values at K=32767). Its benchmark CSV rows had every hardware column blank. Co-Authored-By: Claude <noreply@anthropic.com>
About 1 run in 3 on gfx1250-strict hardware, this config fails validation with wrong values: problem (512,512,1,K) with odd K near 32768 (32767 and 32769), bf16 TN with bias + ReLU, GSU 2-3 with multi-buffer reduction (GSUA MB/MBSK). A different solution fails each time, which points to a race in the GSU multi-buffer reduction rather than one kernel. It failed in 2 of the 3 gfx1250-strict CI runs that reached TensileLite. Add skip-gfx1250-strict until that's fixed, so the per-PR gfx1250 check doesn't go red on unrelated PRs. xfail doesn't fit: with xfail_strict, the runs that pass would fail as XPASS. Co-Authored-By: Claude <noreply@anthropic.com>
talumbau
approved these changes
Oct 13, 2026
talumbau
left a comment
talumbau
left a comment
Collaborator
There was a problem hiding this comment.
Re-reviewed at e3ee5bc. All seven comments are addressed, and the run on this head is green end to end. Approving.
Good catch on the test-selection comment — you were right. Tests/common/config_helpers.py:164-166 turns skip-<arch> into a real pytest.mark.skip for every arch in availableArchs at collection time, so dropping the not skip-gfx1250 term was safe on its own. My comment was based on Tests/conftest.py::pytest_collection_modifyitems, which only auto-skips comm, and I didn't go further. Your choice of not skip-gfx1250-strict over not skip-gfx1250 is also the right one, since base-only skips should still run on strict.
The rest looks right to me: the $(dirname "$LIB") fix, the !cancelled() && steps.strict.outcome == 'success' gate, the read-only cache mount with the prune confined to the nightly, and the hdr = NF / if (NF < hdr) next guard. The awk subscript math holds for every row now — fe keys run 0..NF_hdr-1, so NF >= hdr makes NF - k at least 1. Running TensileLite after a hipBLASLt build failure also seems reasonable given tox builds its own client and tox.ini already carries cmake and invoke in deps.
Two small things, neither blocking:
-
cmakeis still missing from the gate regex while it's still in the sparse checkout.shared/rocrollerwas correctly dropped from both, so the two lists agree now except for this one entry — a PR touching only top-levelcmake/skips the job, and the break surfaces in the nightly after merge. -
Should
ROCR_VISIBLE_DEVICESbe required? Dropping therocminfofallback means the job now hard-fails on any runner whose.envdoesn't set it. That's defensible — it's a clear error instead of a silent multi-GPU run — but it does make the workflow depend on runner configuration that isn't visible from this repo. Was that the intent, or would you rather keeprocminfoas a fallback when the variable is unset?
Contributor
Author
|
Thanks TJ!
Making this required matters because |
1 task done
1 task done
jichangjichang
pushed a commit
that referenced
this pull request
Oct 15, 2026
…formerEngine (#13300) FP8, MXFP8, GELU_AUX support JIRA ID : https://amd-hub.atlassian.net/browse/ROCM-32461 MXBF8 support JIRA ID : https://amd-hub.atlassian.net/browse/AIHPBLAS-5158 ## Motivation TransformerEngine core GEMM tests on gfx1250 fail with "no algorithm" because hipBLASLt has no matching kernel for: - bf16/fp32 `GELU_AUX(_BIAS)` (only fp16 Aux ships); - per-tensor FP8 -> FP8/BF8, because TE always requests AmaxD and no gfx1250 kernel outputs it; - MXFP8 with mixed e4m3/e5m2 inputs, or fp8 x fp8 -> fp16 (also used by DqGEMM). gfx1250 also has no MX kernel with bf8 inputs, so the bf8 rows of `matmul_fast_check_mx_fp8{,_large_k}` (#13031) fail with "NO solution found" and were limited to gfx950 by #12983 (AIHPBLAS-5158). <!-- Explain the purpose of this PR and the goals it aims to achieve. --> ## Technical Details New GridBased logic for gfx1250 and gfx1250-strict: - **Aux:** BF16 and FP32 GEMM + Bias + Act with Aux output (`UseE`), all 4 transposes. Uses the existing non-Aux BF16/FP32 solutions. - **FP8 AmaxD:** per-tensor {F8, B8, F8B8, B8F8} -> {F8, B8} with `OutputAmaxD`, all 4 transposes. Uses the existing ScaleCD solutions. - **MXFP8 (TN):** F8xB8 and B8xF8 -> BF16/F16/F32, F8xF8 -> F16, and B8xB8 -> F32. Uses the existing MX solution. All new solutions build exactly the kernels of the solutions they are derived from. Logic was generated from Tensile configs with `LibraryLogic`. Fixes: - **`PackData_BF16` (tensilelite):** on `HasBF16CVT` archs, the BF16 Aux store converted its own destination register instead of the input values. This produced a garbage Aux output on full tiles with gwvw > 1. It also changes 2 of the existing gfx950 BBS Aux kernels. - **Test client:** for fp8 inputs with an fp8/bf8/f16 output, a bf16 bias was put in an f32 buffer while the descriptor declared bf16. New gtests (`gpu_arch: 1250`): - `matmul_bias_gelu_aux_bf16_fp32{,_large}` - `matmul_fp8_dst_fp8_amaxD_1250` - `matmul_mxfp8_mix_dst_f16_bf16_f32_1250` The bf8 rows of `matmul_fast_check_mx_fp8{,_large_k}` run on `(950|1250)` again. <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan - hipblaslt-bench on TE's shapes (768x4096x768, 2304x4096x768, 256^3) and hipblaslt-test (new gtests + existing `*matmul_real_1b_dst_f8_SCD*`, `*1250_smoke*`, `*matmul_fast_check_mx_fp8*`), each on both library subtrees (gfx1250 and gfx1250-strict code objects) - Functional validation on the gfx1250 FFM and B0 <!-- Explain any relevant testing done to verify this PR. --> ## Test Result FFM, both subtrees (gfx1250 and gfx1250-strict): - Tensile: every new logic file passes validation (D, Aux and AmaxD outputs) — bf16/fp32 Aux, FP8 AmaxD, mixed MXFP8, BF8xBF8 MX. - hipblaslt-bench: all TE shapes pass — bf16/fp32 GELU_AUX_BIAS (4 transposes), FP8/BF8 output with AmaxD (all 8 input/output combinations, plus fp16 and bf16 bias), MXFP8 mixed and ->f16, DqGEMM 256^3. Every case selects the new kernel (previously 0 solutions). - hipblaslt-test: `*matmul_bias_gelu_aux*` 392/392, new FP8 gtests 440/440, regression set 4664/4664 — each on both subtrees. - hipblaslt-test: `*matmul_fast_check_mx_fp8*-*placement*` 12/12 on each subtree, including the 8 bf8 tests that failed with "NO solution found" before. <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: mengzcai <mengzcai@amd.com>
Merged
2 tasks done
Contributor
TheRock Submodule Bump ActivityNewest first
|
tony-davis
pushed a commit
that referenced
this pull request
Oct 20, 2026
JIRA ID : N/A ## Motivation Email a test report after every gfx1250-strict nightly. **NEW:** Report only, per review; the #12983 follow-ups moved to #13771. ## Technical Details - **Report step:** the test job's last step before cleanup runs the new `.github/scripts/blas_nightly_report.py` (standard library only) and emails through a Power Automate flow (secret `GFX1250_BLAS_NIGHTLY_WEBHOOK_URL`). - Sends on the nightly however the tests ended, and on a manual run with `send_report`; never on PRs. A step, so PRs show no skipped check. No runner, no email; a failed send doesn't fail the job. - Step results come from the workflow (`steps.<id>.outcome`, `job.status`); per-test results from the test logs. - Email: banner, run metadata, per-suite counts, and a per-test comparison with the latest nightly in 7 days (new / existing / fixed / skipped or removed). Test lists at the end under Failures, tagged new / timeout / crashed. **NEW:** Lists read `hipBLASLt:` / `(failing)`, one test per line; the comparison's age counts from this run's start, not from when the email is sent. - JSON for the flow: `outcome` (`passed` / `setup_failed` / `tests_failed` / `tests_incomplete` / `cancelled`), `failed_step`, `steps`, per-test `status` (`failing` / `timeout` / `crashed`). - **Workflow:** `send_report` input, `actions: read` (previous nightly's results), `.github/scripts` in the sparse checkout, step ids. ## Test Plan - Real nightly and PR artifacts, plus crafted logs: alarm-timeout shard, crash mid-test, pytest timeout, no GPU, cancelled job, no previous nightly. - Manual run of this branch with `send_report=true`. - Failed send: posted to an unused local port, so nothing is emailed. - `--steps` as the workflow builds it, including steps that never started. - **NEW:** A log whose first pytest session dies before its summary: the later session keeps its own results (the old parser reported them under the first suite). ## Test Result - Email delivered (HTTP 202) from https://github.com/ROCm/rocm-libraries/actions/runs/38626703691: ~68k hipBLASLt, 111 TensileLite common, 264 rocisa passed. - Every crafted case gets the right `outcome`, `failed_step` and per-test status, e.g. "5 failing tests (2 timeout, 1 crashed)". - Failed send logs only `Report not sent (ConnectionRefusedError)`, no URL. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low: CI-only. The step runs only on the nightly or a manual run with `send_report`. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com>
talumbau
pushed a commit
that referenced
this pull request
Oct 22, 2026
…ite timeout; cleanup (#13771) JIRA ID : N/A ## Motivation Follow-ups from the #12983 and #13464 reviews, split out of #13620, plus a fix for hipBLASLt tests running ~7x slower since per-PR runs started. ## Technical Details - **hipBLASLt test speed:** `OMP_NUM_THREADS` and `OPENBLAS_NUM_THREADS` = host CPUs / (4 runners x 4 shards); `BLIS_NUM_THREADS` = the same, at most 8 (hipBLASLt's own BLIS default, which setting `OMP_NUM_THREADS` turns off). By default the CPU reference's OpenMP and OpenBLAS pools each start one thread per host CPU, so 4 runners x 4 shards oversubscribe the host 16x per pool: since per-PR runs started (10-12 to 10-15), the step took ~80 min when other runners were testing vs ~15 min alone, up to 110 min of its 120-min timeout. - **Gate:** add `cmake/` (rocisa and stinkytofu include `cmake/modules/default_amdclang.cmake`). - **One-GPU check:** runs first in its step, under `set -euo pipefail`. `ROCR_VISIBLE_DEVICES` comes from the runner's `.env`; if it is unset, N comes from the runner name's `-gpuN` (the pool registers runner N as `...-gpuN` with `ROCR_VISIBLE_DEVICES=N`), with a warning. It must be a single GPU index (`^[0-9]+$`), or the job fails before the build. Drop `HIP_VISIBLE_DEVICES` (never forwarded into the container). - SDK pinned to `10.2.0a20261013` (compiler `c7ceccf9a5`). Per-PR CI has since moved to compiler `bdc35841c9` (#13835, 10-19), so the workflow's compiler check now warns; re-pinning to `10.2.0a20261018` or later, which carries `bdc35841c9`, is a follow-up. - **upload-artifact:** v4.6.2 -> v7.0.2 (Node 20 deprecated). v7 needs Actions Runner >= 2.327.1; the gfx1250 runner is 2.337.0. - **TensileLite per-test timeout:** 600 -> 1200 s, this workflow only (FFM and TheRock keep their own). `decouple_pgr_tdm_fuse_gfx1250.yaml` has timed out at 600 s in most runs since #12825 grew its sweep 4.5x; before, it passed 41 of 41. - **Header:** the workflow is marked temporary, to be removed once TheRock's multi-arch builds produce gfx1250-strict artifacts (ROCm/TheRock#8350) and gfx1250-strict runners are available in TheRock CI. ## Test Plan - OpenMP cap: 4 jobs x 4 shards of a `*pre_checkin*` slice on a 128-core gfx942 host, default vs capped. - One-GPU check: 7 values under `bash -euo pipefail`. ## Test Result - [x] Run https://github.com/ROCm/rocm-libraries/actions/runs/38994215757: all steps green with no warnings on SDK 10.2.0a20261013 - hipBLASLt 68,925 tests, 0 failed, in 3.9 min (OpenMP/OpenBLAS capped, BLIS at most 8) while an uncapped run on the same host took 52 min - TensileLite 113 passed, including decouple_pgr_tdm_fuse_gfx1250.yaml under the 1200 s timeout. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low: CI-only, gfx1250 workflow. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com>
jinchen62
added a commit
that referenced
this pull request
Oct 23, 2026
…rict The strict gfx1250 runner reports gfx1250-strict, which skip-gfx1250 does not match, so the per-PR gfx1250-strict job built these gfx950-only configs and failed. Add skip-gfx1250-strict, as every other gfx950-only config has since #12983. Co-authored-by: Cursor <cursoragent@cursor.com>
tony-davis
added a commit
that referenced
this pull request
Oct 23, 2026
…trict (#13504) JIRA ID : AIHPBLAS-5134 ## Summary About two thirds of hipBLASLt's architecture-filtered client test cases never ran on gfx1250: 132 of 196 entries in `matmul_gtest.yaml`, `smoke_gtest.yaml` and `rocroller_gtest.yaml`. Most were scoped to gfx942 or gfx950, including the fast_check address-arithmetic cases. This PR reviews each one against what the gfx1250 libraries actually provide, and enables those that apply. - **Enabled on gfx1250 (36 entries):** 32 cases get `1250` added to their `gpu_arch` pattern, and 4 grouped-GEMM cases get a corrected exclusion. Each was checked type by type and layout by layout against the gfx1250 logic files (`library/.../Logic/asm_full/gfx1250/`), so every enabled combination has kernels there. - fast_check: `size_t_index_overflow`, `size_t_output_overflow_epilogue`, `offset_above_2_32`, `batch_offset_above_2_32`, `store_extent_ROCM31016`, `size_t_output_overflow_all_solutions`, `every_solution_repeat_xf32`, `large_k_fp8`, `mixed_fp8`, the five FP8 scale cases, `matches_unit_check_bias_grad_bf16`, `threshold_custom_kernel_ROCM31258`. - Other address-arithmetic cases: `size_t_index_overflow{,_hpa}`, `size_t_output_overflow_bf16_ROCM26455`, `size_t_batch_stride_overflow`, `integer_exact_custom_kernel_store_extent_xf32_ROCM31258`, `batch_offset_large`. - Data types and features: `gemm_double` (and its smoke case), `gemm_i8_dst_i32`, `gemm_i8_dst_i8`, `gemm_i8_dst_i32_alpha_float_bits`, `ext_i8_dst_i32_support_smoke`, `bgradb_bf16_NT`, `stride_lt_ld`, `multithread_stream`, `extapi_algo_method_tuning_gsu_gemm_only`. - Corrected exclusion: `groupedgemm_zero_n`, `extapi_groupedgemm` (2), `extapi_algo_method_grouped_gemm_tuning_wgm`. Their `gpu_arch_exclude: '1[1-2]\d{2}'` was meant for consumer gfx11/gfx12 but also matched gfx1250. It's now `'11\d\d|12[^5]\d'`, the form other cases in the file already use. - **Kept off gfx1250 (96 entries), each for a reason:** - **fnuz FP8 (42).** gfx1250 has OCP FP8 only. - **Architecture-specific cases (36):** consumer gfx11/gfx12 grid-limit, bias and i8 cases; gfx90a compute variants; gfx942 swizzle, conv3d and `_94x`/`_gfx942` variants; gfx950 subtile FP4, tf32x1 and custom-kernel cases; rocRoller (gfx1250 builds skip it). - **No gfx1250 kernels for the problem type (18).** Grouped GEMM for bf16 and xf32; fp32 and bf16 aux-output (E) epilogues; the fp32 NN stagger-increment overflow case (`size_t_stagger_incs_overflow_f32_ROCM31230`, which the first gfx1250-strict run found has no solution); fp32 and f32-output bias gradients; MX cases using FP6 or fp16 output; f32 with fast bf16 compute in TT. These, and the gfx950 custom-kernel case, carry a `# Not on gfx1250:` comment. A case with no solution fails rather than skips, so enabling them would only add red. - **gfx1250-strict** runs the same cases: every pattern containing `1250` matches both device names. Strict-only failures are quarantined with `known_bug_platforms: "gfx1250-strict"`, which matches exactly. `TESTING.md` documents this, plus the consumer-exclusion idiom. - **Harness:** `gpu_arch` and `gpu_arch_exclude` widen from 16 to 32 bytes, since `9(0a|42|50)|1250` doesn't fit in 16. `hipblaslt_gentest.py` now refuses any char-array value that doesn't fit with its terminator. Before, ctypes silently stored a value that exactly filled the field with no NUL. No library code changes. ## Work tracking - [x] Internal ticket: [AIHPBLAS-5134](https://amd-hub.atlassian.net/browse/AIHPBLAS-5134) (sub-task of AIHPBLAS-4988) - [ ] GitHub issue: - [ ] No tracker: reason: ## Tests - [x] New work: tests cover the change - [ ] Defect fix: regression test included - [ ] Defect: test-only PR: known-bugs added (W-KNOWN-BUG); tracker in YAML - [ ] Fix PR: known-bugs entry removed: ______ - [ ] No test changes: waiver: ______ **Paths:** - `projects/hipblaslt/clients/tests/data/matmul_gtest.yaml`, `smoke_gtest.yaml`: arch patterns and `# Not on gfx1250:` comments - `projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml`, `clients/common/include/hipblaslt_arguments.hpp`: field widths - `projects/hipblaslt/clients/tests/hipblaslt_gentest.py`, `clients/tests/test_hipblaslt_gentest.py`: length check and its tests - `projects/hipblaslt/TESTING.md`: "Scoping test cases to architectures" **Test results:** - Test data generation: `hipblaslt_gentest.py data/hipblaslt_gtest.yaml` succeeds. Output grows by exactly 32 bytes per record over 119,020 records, so the case count is unchanged. No existing char-array value reaches its limit. - `python -m unittest discover -s clients/tests -p 'test_hipblaslt_gentest.py'`: 4 pass. With the check weakened, 2 fail. These run by hand, like `clients/scripts/sdc_hunt`; no CI job runs client Python tests today. The check itself runs in every build's test data generation. - gfx942/gfx950: each edit only appends `|1250` or narrows an exclusion away from gfx1250, so the case set there is unchanged. Precheckin on this PR: pending. - gfx1250-strict: the first `gfx1250 HW hipBLASLt + TensileLite Tests` run on this PR passed 17,373 hipBLASLt tests and failed 1: `size_t_stagger_incs_overflow_f32_ROCM31230` ("NO solution found!"). That enable is reverted in c08a775 with the reason recorded. A rerun on that commit is pending. The run's one TensileLite failure was a 600 s timeout in `decouple_pgr_tdm_fuse_gfx1250.yaml`, which this PR doesn't touch and nightly develop runs have passed. ## Flags / guardrails - [x] N/A. Test data and test-harness changes only. - [ ] Guarded: name: ______ default: ______ - [ ] Unguarded: waiver: ______ approver: ______ ## Adjacent tests considered - TensileLite pytest suites have their own gfx1250 skip and xfail gating (AIHPBLAS-4764). That's reviewed separately, not here. - `matmul_fast_check_mx_*` already runs on gfx1250 and is the MX coverage there, so the `matmul_mx_*` cases that need FP6 or fp16-output MX stay off. #13464 narrows the bf8 rows inside those fast_check MX cases for the same reason; it touches different entries from this PR. - #12983 adds a per-PR gfx1250-strict hardware job. Once it lands, this PR's enabled cases run there automatically. - The `_94x` i8 variants add activation epilogues. They stay off until i8 activation support on gfx1250 is confirmed; the plain i8 cases are enabled. ## Risk acceptance - [x] STANDARD - [ ] WAIVER: code: ______ approver: ______ follow-up: ______ [AIHPBLAS-5134]: https://amd-hub.atlassian.net/browse/AIHPBLAS-5134?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ --------- Co-authored-by: Cursor <cursoragent@cursor.com>
ScottTodd
requested changes
Oct 27, 2026
Comment on lines
+226
to
+233
| REF="$(sed -nE '/therock-ref:/,/value:/ s/.*value: *"([0-9a-f]{40})".*/\1/p' .github/actions/ci-env/action.yml)" | ||
| export REF | ||
| python3 - <<'EOF' | ||
| import glob, json, os, sys, urllib.request | ||
|
|
||
| # Warn-only: compare the pinned SDK's compiler with the one per-PR CI builds with (the | ||
| # amd-llvm at ci-env's therock-ref). The pin stays fixed; a mismatch is reported, never fatal. | ||
| ref, sdk = os.environ.get("REF", ""), os.environ.get("ROCM_SDK_VERSION", "") |
Member
There was a problem hiding this comment.
Inline code in workflow files like this is totally unacceptable and must be reviewed by CI/CD maintainers or release admins (cc @marbre ). Each instance using ci-env/action.yml and therock-ref like this puts our continuous releases at risk.
Workflows should follow https://github.com/ROCm/TheRock/blob/main/docs/development/style_guides/github_actions_style_guide.md
1 task done
ScottTodd
added a commit
that referenced
this pull request
Oct 28, 2026
## Motivation * See ROCm/rockrel#88 ## Technical Details * #14007 was included on the release branch, so fewer files need updating Generated with Codex. This *should* be deterministically automatable, but... * `.github/actions/ci-env/action.yml` still has uses when it shouldn't, see #12720 (comment) * `.github/workflows/gfx1250_hipblaslt_tensilelite_test.yml` is new and incompatible with release branches, see #12983 (review) ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : N/A
Motivation
Technical Details
rocm_sdk_version, default10.2.0a20261007). A warn-only step flags when the SDK's compiler differs from the one at ci-env'stherock-ref.gfx1250-strict(HSA_DISABLE_GFX12_STRICT=0), then builds withinstall.sh -a gfx1250-strict --skip_rocroller.ROCR_VISIBLE_DEVICES(-e ROCR_VISIBLE_DEVICES) and fails early unless it names exactly one GPU. Checked on the host during a CI run: all 4 hipBLASLt shards ran on the runner's own GPU.library/gfx1250-strict/; fails ifhipblaslt-testdoesn't link the library built here or its code objects aren't labelledgfx1250-strict. Runs as 4 gtest shards.-m "common and not skip-gfx1250-strict"; skipped if setup failed, still runs after a hipBLASLt failure.pull_requesttrigger behind a GitHub-hosted gate (non-draft, ROCm org label, changes underprojects/hipblaslt, its shared deps, this workflow orci-env/action.yml). One run per PR; a newer push cancels it. Manual dispatch and the nightly are ungated. A manual run and the nightly on the same branch don't cancel each other.actions/setup-python, so nothing from the runner's tool cache); the runner's pip cache is used only by the SDK install and is read-only except for the nightly; the job hands the root-owned workspace back to the runner's user at the end.skip-gfx1250-stricton 4 TensileLite configs and the bf8 MX gtest rows limited to gfx950, + fixed comment saying gfx1250 has no MX kernels with bf8 inputs.gsu_gfx1250-strict.yamlon gfx1250-strict for now (skip-gfx1250-strict): it intermittently returns wrong values from GSU multi-buffer kernels (see Test Result). Followed up with the config owner; the skip comes out when that's fixed.workflow_runtrigger, artifact resolution, commit-status posting, unusedGFX1250_HW_XFAIL.Test Plan
gfx1250-strict hipBLASLt + TensileLitecheck on the gfx1250 runners (the PR changes the workflow, so it runs here).Test Result
Submission Checklist
Risk level
🤖 Generated with Claude Code