Skip to content

fix(hipblaslt): pin gfx942 build-only test to its target arch - #12086

Merged
talumbau merged 1 commit into
developfrom
users/talumbau/fix-gfx942-build-only-gpu-targets
Sep 16, 2026
Merged

talumbau merged 1 commit into
developfrom
users/talumbau/fix-gfx942-build-only-gpu-targets

Conversation

@talumbau

@talumbau talumbau commented Sep 15, 2026 •

Copy link
Copy Markdown
Collaborator

JIRA ID : AIHPBLAS-4748

Motivation

Fixes #12083. Tensile/Tests/common/flags/gfx942/test_build_only.py::test_compile fails on gfx1250v0-revision hosts with AssertionError: Unhandled WMMA: (16, 16, 16, 1), even though its embedded config hardcodes ISA: [[9,4,2]] (gfx942) and the test is documented to run on any machine as a hardware-independent, --build-only compile check.

Root cause: tensile_args (Tensile/Tests/conftest.py) forwards pytest's --gpu-targets option into every test's Tensile.Tensile(...) invocation when the harness sets one (e.g. tox.ini injects --gpu-targets gfx1250v0 whenever invoke get-gpu-revision-target reports gfx1250v0). Tensile.py gives CLI --gpu-targets priority over the config's own ISA: field, so the test's gfx942 target was silently overridden to gfx1250v0. gfx1250's WMMA v3 ISA only supports K-depths of 4/32/64/128 (Tensile/Common/Utilities.py::wmmaV3InputVgprLayout), so the config's (16, 16, 16, 1) MFMA-only instruction shape (valid on gfx942) hit the assert False, f"Unhandled WMMA: {wmma}" branch — a false-positive CI/local-test failure unrelated to the actual gfx1250 kernel work being tested.

Technical Details

Pin --gpu-targets gfx942 explicitly in both test_compile and test_use_cache, appended after *tensile_args in the argument list passed to Tensile.Tensile(...). Tensile.py's --gpu-targets argparse option uses the default store action, so the last occurrence on the command line wins — this guarantees the test always builds for gfx942 regardless of what the harness forwards, without touching the legitimate --gpu-targets forwarding behavior other tests rely on.

test_use_cache was not actually exposed to this bug today (it's guarded by its own skipif(not has_arch("gfx942")), so it only ever runs on real gfx942 hardware where any forwarded --gpu-targets would already say gfx942), but it makes the same "always gfx942" claim as test_compile, so it gets the same explicit pin for consistency and to guard against future changes to the skip condition or harness behavior.

Test Plan

  • Reproduced the failure on develop via the standard CI flow: tox -e py3 -- Tensile/Tests -m common --durations=0 on a 4-GPU gfx1250/gfx1250v0 host, both with the default 4 pytest workers and with TENSILE_NUM_PYTEST_WORKERS=1.
  • After the fix, ran the specific test directly against the prebuilt client, reproducing the same --gpu-targets gfx1250v0 forwarding tox performs:
    pytest --prebuilt-client=... --global-parameters CheckASMCodeSize=True --tensile-options=--gpu-targets,gfx1250v0 Tensile/Tests/common/flags/gfx942/test_build_only.py
  • Spot-checked an unrelated gfx1250 config test (gemm/gfx12/tdmpluslds_tdm_gfx1250.yaml) under the same --gpu-targets gfx1250v0 forwarding to confirm the fix doesn't affect tests that rely on that forwarding.

Test Result

Before fix:

FAILED Tensile/Tests/common/flags/gfx942/test_build_only.py::test_compile - AssertionError: Unhandled WMMA: (16, 16, 16, 1)

After fix:

Tensile/Tests/common/flags/gfx942/test_build_only.py::test_compile PASSED
Tensile/Tests/common/flags/gfx942/test_build_only.py::test_use_cache SKIPPED (gfx942 GPU not available)

Unrelated gfx1250 config test, same environment, still passes:

Tensile/Tests/common/test_config.py::test_config[Tensile/Tests/common/gemm/gfx12/tdmpluslds_tdm_gfx1250.yaml] PASSED

Submission Checklist

Risk level

Low. The change is scoped to two test functions in a single test file, only adds an explicit --gpu-targets gfx942 argument (highest-precedence override) to calls that were previously letting that value fall through from the harness; it does not change any non-test code path or other tests' behavior.

…arch

test_compile (and test_use_cache) in test_build_only.py hardcode ISA
gfx942 in their embedded config and are meant to build regardless of
host architecture, but tensile_args forwards any --gpu-targets set by
the harness (e.g. tox injects --gpu-targets gfx1250v0 on gfx1250v0
hosts), which takes priority over the config's own ISA in Tensile.py.
On such hosts the test silently compiled for gfx1250v0 instead of
gfx942, and gfx1250's WMMA v3 ISA doesn't support the config's
(16,16,16,1) MFMA-only instruction shape, causing an assertion
failure unrelated to the actual gfx1250 kernel work under test.

Pin --gpu-targets gfx942 explicitly after *tensile_args so it always
wins, regardless of what the harness forwards.

Fixes #12083

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
@therock-pr-bot

therock-pr-bot Bot commented Sep 15, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass PR does not contain code files — Unit Test auto-passed
🔎 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

Copy link
Copy Markdown

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

@codecov

codecov Bot commented Sep 15, 2026 •

Copy link
Copy Markdown

Codecov Report

✅ All modified and coverable lines are covered by tests.

❌ Your project check has failed because the head coverage (46.40%) 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   #12086      +/-   ##
===========================================
- Coverage    70.55%   70.55%   -0.00%     
===========================================
  Files         2810     2810              
  Lines       462535   462575      +40     
  Branches     68105    68113       +8     
===========================================
+ Hits        326312   326328      +16     
- Misses      112680   112699      +19     
- Partials     23543    23548       +5     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 46.40% <ø> (ø)
TensileLite-Unit 76.08% <ø> (+0.01%) ⬆️
hipBLAS 90.62% <ø> (ø) Carriedforward from 46670b6
hipBLASLt 35.24% <ø> (-0.03%) ⬇️
hipCUB 82.68% <ø> (ø) Carriedforward from 46670b6
hipDNN 87.02% <ø> (ø) Carriedforward from 46670b6
hipFFT 42.66% <ø> (ø) Carriedforward from 46670b6
hipRAND 76.12% <ø> (ø) Carriedforward from 46670b6
hipSOLVER 68.92% <ø> (ø) Carriedforward from 46670b6
hipSPARSE 86.99% <ø> (ø) Carriedforward from 46670b6
rocBLAS 48.31% <ø> (ø) Carriedforward from 46670b6
rocFFT 51.72% <ø> (ø) Carriedforward from 46670b6
rocRAND 56.90% <ø> (ø) Carriedforward from 46670b6
rocSOLVER 76.83% <ø> (ø) Carriedforward from 46670b6
rocSPARSE 74.61% <ø> (ø) Carriedforward from 46670b6
rocThrust 91.60% <ø> (ø) Carriedforward from 46670b6

*This pull request uses carry forward flags. Click here to find out more.
see 6 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.

@tony-davis tony-davis left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nice fix, and a genuinely tricky one to track down. I checked the --gpu-targets precedence claim directly in Tensile.py: it's a store action, so the flag you appended after *tensile_args does win over anything the harness forwards. That matches the behavior you described.

A couple of things I dug into while reviewing:

  • I looked at the currently-failing checks on this PR. FFM gfx1250 / HW gfx1250 are the gfx1250 bring-up lane, not required, so I didn't weigh those.
  • mci/rocm-libraries/precheckin(hipblaslt) is also red right now, but I pulled the actual Jenkins log for the failing stage (Test ubuntu24-therock && gfx950). It's a GPU memory fault (HSA_STATUS_ERROR_MEMORY_FAULT, core dump) inside hipblaslt-test on gfx950, in an unrelated matmul_test case. Nothing to do with this diff, so I'm treating it as a pre-existing gfx950 hardware flake.

Also spot-checked that your fix doesn't break legitimate --gpu-targets forwarding elsewhere (the tdmpluslds_tdm_gfx1250.yaml check in your test plan). Good catch adding that.

One tiny, non-blocking thought for later: the docstring explains why the pin is needed but not that --gpu-targets is last-wins on the CLI, which is the fact that makes this fix reliable rather than incidental. Not asking for a change, just flagging it in case it's useful context for the next person who touches this file.

@archana-ramalingam

Copy link
Copy Markdown
Contributor

LGTM. Failing FFM and HW tests are unrelated.

@talumbau
talumbau merged commit 61d8786 into develop Sep 16, 2026
123 of 137 checks passed
@talumbau
talumbau deleted the users/talumbau/fix-gfx942-build-only-gpu-targets branch September 16, 2026 23:46
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

tony-davis added a commit that referenced this pull request Sep 17, 2026
JIRA ID : AIHPBLAS-4510

## Motivation
- 4 gfx1250 TensileLite `common` configs fail on the current gfx1250 HW.
- xfail them so the per-PR and nightly gfx1250 lanes stay green while
the gaps
  are tracked.
  - The separate gfx942 `test_compile` failure is fixed in #12086 

## Technical Details
  - Mark the currently-failing gfx1250 configs `xfail-gfx1250`
- Fix `findAvailableArchs` (`config_helpers.py`) to strip a trailing
version suffix arch — without this the marks never fire on a versioned
runner and the configs stay red.
- FFM path uses HSA_MODEL_MEMFILE to skip any HW specific
`xfail-<arch>`, as these tests failing on HW pass on FFM.
- Add 3 tests in `Tensile/Tests/common/test_config_helpers.py` covering
the arch-suffix strip and the FFM-gate suppression.

## Test Plan
  - 3 new common tests in `test_config_helpers.py`
  - FFM tests pass as expected
- Full gfx1250 HW xfail behavior is exercised once #12164 's workflow
runs post-merge against a rebuilt artifact.

## Test Result
- [x] `pytest Tensile/Tests/common/test_config_helpers.py` → 7 passed (4
existing + 3 new).
- [x] FFM tests pass as expected -
https://github.com/ROCm/rocm-libraries/pull/12178/checks?check_run_id=105181618286
  - gfx1250 HW result pending the post-merge #12164 run.
  
## Submission Checklist

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

## Risk level
- Low — only adds xfail marks to already-failing gfx1250 configs plus a
one-line
arch-normalization with regression tests; no change to any passing test.
`xfail_strict` turns an unexpected pass red, so a fixed config can't
stay silently masked.

---------

Co-authored-by: Tony Davis <tony.davis@amd.com>
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.

fix(hipblaslt/tensilelite): gfx942 build-only test fails on gfx1250v0 hosts due to --gpu-targets override

3 participants