Repository navigation
test(hipblaslt): skip gfx1250-strict where the tests can't run - #13464
Conversation
Four TensileLite configs skip gfx1250 but not gfx1250-strict, and a base
skip-gfx1250 mark does not fire on gfx1250-strict by design, so they run there
and fail before reaching the GPU:
- custom_rocblas_gemv_bf16_tn_m{1,2,4}: gfx950 hand-written assembly
(v_dot2c_f32_bf16, s_waitcnt) that doesn't assemble for gfx12.
- sk_queue_selfclean: SGEMM with a 32x32x2 MFMA tile, rejected as an invalid
WMMA configuration (0 solutions).
Add skip-gfx1250-strict to each.
The MX fast_check tests (#13031) run every fp8/bf8 pair on '(950|1250)', which
also matches gfx1250-strict. gfx1250 and gfx1250-strict have MX kernels for
f8 x f8 only, so the bf8 pairs fail with "NO solution found". Keep f8 x f8 on
'(950|1250)' and run the bf8 pairs on '950', like the other MX tests. The entry
names are unchanged, so gfx950 test IDs stay the same.
Co-Authored-By: Claude <noreply@anthropic.com>
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## users/aramalin/gfx1250-hw-test-runtime-libpath #13464 +/- ##
=================================================================================
Coverage ? 71.92%
=================================================================================
Files ? 2876
Lines ? 475981
Branches ? 70364
=================================================================================
Hits ? 342329
Misses ? 109881
Partials ? 23771
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
Instead of splitting matmul_fast_check_mx_fp8 and _large_k into separate gfx950 entries, tag the four rows with bf8 inputs with gpu_arch '950'. The generator merges each row over its entry, so the generated tests are the same, and the diff drops from 45 changed lines to 13. Co-Authored-By: Claude <noreply@anthropic.com>
|
Failing tests are unrelated existing failures/flaky tests. |
tony-davis
left a comment
There was a problem hiding this comment.
Hi Archana! Thanks for chasing these down. I like the per-row gpu_arch version a lot better than the split.
I checked a few things locally. On both develop and the stack base, these 4 TensileLite configs are the only non-strict configs carrying skip-gfx1250 without skip-gfx1250-strict, so this closes the whole gap (not just the ones CI happened to hit). I also ran hipblaslt_gentest.py on the base and both commits. Base to head changes exactly 8 of 119,022 test records, and the only field that changes is gpu_arch ((950|1250) to 950). Both commits generate the same set of tests, just in a different order. So gfx950 keeps all its MX coverage. The red ASAN and TheRock checks look unrelated to me too.
A few non-blocking thoughts:
- Is the missing bf8 MX support on gfx1250 intended, or is it on someone's list? This also drops the bf8 rows on base gfx1250, not just strict. If it's a planned gap, a ticket referenced in the YAML comment would help us turn these rows back on when the kernels land.
- The YAML comment says gfx1250 has MX kernels "for f8 x f8 only", but there are f8 x f4 and f4 x f4 MX logic files too. Maybe "gfx1250 has no MX kernels with bf8 inputs"?
- There are a couple of typos in the description ("hipblalst").
Side note, not for this PR: hipblaslt_gentest.py output isn't reproducible run to run (the known_bug_platforms set ordering changes with Python's hash seed). Probably worth its own issue.
LGTM!
|
Thanks for the review @tony-davis, and for checking the generated tests against the base.
Addressing in phase 2 set of changes:
|
37f4155
into
users/aramalin/gfx1250-hw-test-runtime-libpath
JIRA ID : N/A ## Motivation - Run hipBLASLt + TensileLite on the gfx1250-strict path (#11777) for every PR. - Per-PR TheRock builds don't produce gfx1250-strict yet (ROCm/TheRock#8350, draft), so the job builds from source instead of waiting for artifacts. - Add Nightly schedule ## Technical Details - Toolchain: pinned ROCm SDK nightly (`rocm_sdk_version`, default `10.2.0a20261007`). A warn-only step flags when the SDK's compiler differs from the one at ci-env's `therock-ref`. - Checks that the runtime reports `gfx1250-strict` (`HSA_DISABLE_GFX12_STRICT=0`), then builds with `install.sh -a gfx1250-strict --skip_rocroller`. - One runner per GPU: the job container gets the runner's `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. - hipBLASLt: loads `library/gfx1250-strict/`; fails if `hipblaslt-test` doesn't link the library built here or its code objects aren't labelled `gfx1250-strict`. Runs as 4 gtest shards. - TensileLite: Uses 4 workers sharing the GPU. Unit tests dropped (host-only, already run per PR in TheRock CI). Selection `-m "common and not skip-gfx1250-strict"`; skipped if setup failed, still runs after a hipBLASLt failure. - Log scrub on build and test output: hardware details and wall-clock times are removed before anything reaches the console or the uploaded logs. - Per PR: `pull_request` trigger behind a GitHub-hosted gate (non-draft, ROCm org label, changes under `projects/hipblaslt`, its shared deps, this workflow or `ci-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. - Nightly: develop at 03:00 UTC (7 pm PST, 8 pm PDT). - Runner state: a fresh venv per job from the image's Python 3.12 (no `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. - Includes #13464 (merged into this branch): `skip-gfx1250-strict` on 4 TensileLite configs and the bf8 MX gtest rows limited to gfx950, + fixed comment saying gfx1250 has no MX kernels with bf8 inputs. - Skips `gsu_gfx1250-strict.yaml` on 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. - Removed: the `workflow_run` trigger, artifact resolution, commit-status posting, unused `GFX1250_HW_XFAIL`. ## Test Plan - This PR's own `gfx1250-strict hipBLASLt + TensileLite` check on the gfx1250 runners (the PR changes the workflow, so it runs here). ## Test Result - CI on the gfx1250 runners ([run 38319193893](https://github.com/ROCm/rocm-libraries/actions/runs/38319193893)): hipBLASLt 68,138 tests over 4 shards, 0 failed; TensileLite common 109 passed, 2 xfailed, 0 failed. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level - Low. Self-contained workflow. PR code runs on a self-hosted runner only for ROCm org-labelled PRs 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com>
…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>
…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>
JIRA ID : NA
Motivation
developonce ci(gfx1250): self-build gfx1250-strict per PR #12983 merges.Technical Details
TensileLite: add
skip-gfx1250-strictto configs that already skipgfx1250. A baseskip-gfx1250mark 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(feat(hipblaslt): add rocBLAS wvSpltK BF16 TN skinny GEMMs (n ≤ 4) via GFA #12329): gfx950 hand-written assembly (v_dot2c_f32_bf16,s_waitcnt) that doesn't assemble for gfx12.streamk/sk_queue_selfclean.yaml(feat(tensilelite): check the shared Synchronizer buffer is left clean, by default #10451): SGEMM with a 32x32x2 MFMA tile, rejected as an invalid WMMA configuration (0 solutions).hipBLASLt
matmul_gtest.yaml(test(hipblaslt): check MX block scales exactly with fast_check #13031):matmul_fast_check_mx_fp8and_large_krun 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 withNO solution found!. The four rows with bf8 inputs getgpu_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
Test Result
Submission Checklist
Risk level
🤖 Generated with Claude Code