test(tensilelite): add true16 v_max_f16 AMax activation regression#9711
Open
tony-davis wants to merge 5 commits into
Open
test(tensilelite): add true16 v_max_f16 AMax activation regression#9711tony-davis wants to merge 5 commits into
tony-davis wants to merge 5 commits into
Conversation
…OCM-3994) Add a RED-on-develop reproducer for the true16 (real-true16) activation-clamp defect: AMaxKernelGenerator emits the fake16 form of the v_max_f16 activation clamp, which real-true16 assemblers reject. - Primary emit-site test (Tensile/Tests/unit/test_amax_true16_activation.py): drives AMaxKernelGenerator for a half input and asserts max_per_data / merge_sum emit the true16 form (.l selector on every operand, suffix inside the abs() paren). Host-only string assertion, no GPU. Parametrized across the gfx11 family the NoSDWA cap enables (gfx1100/1101/1102/1103/1150/1151/1152/1153). - Secondary rocisa unit test (rocisa/test/test_instruction.py): asserts the EMaxF16 true16 rendering directly. Both are quarantined with a strict xfail (W-KNOWN-BUG two-PR flow): they run and xfail on develop, and strict=True flips the XPASS to a failure once the true16 fix lands, forcing removal of the markers. Test-only; no production source changed. JIRA ID : ROCM-3994 Co-authored-by: Cursor <cursoragent@cursor.com>
tony-davis
changed the title
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Contributor
There was a problem hiding this comment.
Pull request overview
Adds host-only regression coverage for ROCM-3994, pinning the expected true16 encoding of the v_max_f16 activation-clamp emitted by TensileLite/rocisa, while quarantining the tests via strict xfail as part of a two-PR known-bug flow (test lands ahead of the production fix).
Changes:
- Add a new TensileLite unit test that renders
AMaxKernelGeneratoroutput for half inputs and assertsv_max_f16uses true16 half-word selectors (and that the.l/.hsuffix is insideabs(...)). - Add a rocisa unit test that asserts the true16 rendering of the
EMaxF16helper (activation/merge shapes), also quarantined with strictxfail.
Reviewed changes
Copilot reviewed 2 out of 2 changed files in this pull request and generated 1 comment.
| File | Description |
|---|---|
| projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_amax_true16_activation.py | New host-only parametrized regression tests validating true16 v_max_f16 emission in AMaxKernelGenerator (quarantined with strict xfail). |
| projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py | New rocisa-level unit test validating EMaxF16 true16 string rendering (quarantined with strict xfail). |
…test Address PR review: resolve the assembler path with ROCM_PATH + shutil.which (matching test_mubuf.py::_isa_context) instead of the bare "amdclang++" string, so rocIsa.init works when ROCm's bin dir is not on PATH. Falls back to the bare name if not found. JIRA ID : ROCM-3994 Co-authored-by: Cursor <cursoragent@cursor.com>
davidd-amd
reviewed
Jul 22, 2026
davidd-amd
reviewed
Jul 22, 2026
davidd-amd
reviewed
Jul 22, 2026
davidd-amd
approved these changes
Jul 22, 2026
davidd-amd
left a comment
davidd-amd
left a comment
Contributor
There was a problem hiding this comment.
I think that the existence of the test as self documenting code and a improved safety net is the most important thing. The testing has issues from the perspective of unit testing but we can't fix that now and this gets us to a better place in the near term.
This was referenced Jul 22, 2026
…ain decouple Address PR review (davidd-amd): add a brief comment explaining that the test constructs EMaxF16, renders it via str(), and asserts on the emitted assembly form (no assembler/GPU invoked), and note what vgpr(...) constructs. Add a TODO(#9720) to decouple these string-rendering tests from the toolchain/env init. JIRA ID : ROCM-3994 Co-authored-by: Cursor <cursoragent@cursor.com>
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #9711 +/- ##
===========================================
- Coverage 69.50% 68.55% -0.95%
===========================================
Files 2770 2764 -6
Lines 452498 459457 +6959
Branches 66620 71410 +4790
===========================================
+ Hits 314492 314968 +476
- Misses 117526 122096 +4570
- Partials 20480 22393 +1913
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
The true16 AMax activation regression test (ROCM-3994) imports AMaxGenerator directly from the tensilelite root, but that script was never installed into share/hipblaslt/tensilelite/. In the CI artifact tree the bare `import AMaxGenerator` raised ModuleNotFoundError, which aborted the whole Tensile/Tests/unit collection and reddened the tensilelite shard. Ship AMaxGenerator.py at the install root alongside the other test artifacts so the regression test can actually run in CI. Co-authored-by: Cursor <cursoragent@cursor.com>
tony-davis
enabled auto-merge (squash)
July 24, 2026 15:50
Alex-Vasile
approved these changes
Jul 24, 2026
Alex-Vasile
left a comment
Alex-Vasile
left a comment
Contributor
There was a problem hiding this comment.
looks good. Thanks for the test
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 : ROCM-3994
Motivation
hipBLASLt/TensileLite emits the fake16 form of the
v_max_f16activation-clamp instruction inAMaxKernelGenerator. Under a compiler that enforces real16 mode (+real-true16), the assembler rejects it (operands are not valid for this GPU or mode). PR #6652 migrated the conversion instructions to true16 but not this activation path, so the defect escaped with no test pinning it. This PR supplies the missing regression test ahead of the fix.Technical Details
The correct true16 form binds a half-word selector to each register operand, with the suffix placed inside the
abs(...)paren:v_max_f16 v[out].l, v[out].l, abs(v[val].l).Two host-only tests (no GPU; the target arch is forced from a gfx string):
Tensile/Tests/unit/test_amax_true16_activation.py): constructs anAMaxKernelGeneratorfor a half input, rendersmax_per_data/merge_sumto a string, asserts the true16 half-word form and rejects the bare fake16 form. Parametrized across the gfx11 family enabled by theNoSDWAcap (gfx1100/1101/1102/1103/1150/1151/1152/1153) rather than a single arch.rocisa/test/test_instruction.py::test_instruction_vmax_f16_true16): asserts theEMaxF16true16 rendering directly, in the style of the existingtest_instruction_cvt.Both are quarantined with a strict
xfailunder the W-KNOWN-BUG two-PR flow, tracked by ROCM-3994 and time-boxed to the next hipBLASLt release. They still execute every CI run; on develop they xfail, andstrict=Trueturns the XPASS into a hard failure the moment the true16 fix lands, forcing the fix PR to delete the markers. This is a deliberate quarantine, not a disabled test. Test-only change; no production source modified.Test Plan
Run the two tests on
develop(bug present) and on the true16 fix branch (users/ericwan/true16-patch, after rebuilding rocisa):pytest rocisa/test/test_instruction.py -k vmax_f16_true16 -qpytest Tensile/Tests/unit/test_amax_true16_activation.py -qCI: the secondary runs in the rocISA component lane (Linux gfx94X + Windows gfx1151); the primary runs in the TheRock
blasmatrix'stensilelitetest step.Test Result
17 xfailed(1 secondary + 16 primary), suite green. Failure reasons confirmed: primary emitsv_max_f16 v[vgprOutput], v[vgprOutput], v[vgprOutputB](fake16); secondary fails on the absentEMaxF16import.16 passed, secondary1 passed; emitted formv_max_f16 v[vgprOutput].l, v[vgprOutput].l, abs(v[vgprValue+0].l). With the strict-xfail markers in place, this becomes an enforced XPASS->fail that the fix PR resolves by removing them.Submission Checklist
Risk level
None. Test-only, host-only, quarantined; no production source or goldens touched.
Made with Cursor