From b4545af3b73e8c15b560bc76b14a50fce1c34665 Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Wed, 22 Jul 2026 12:41:00 -0400 Subject: [PATCH 1/4] test(tensilelite): add true16 v_max_f16 AMax activation regression (ROCM-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 --- .../Tests/unit/test_amax_true16_activation.py | 176 ++++++++++++++++++ .../rocisa/test/test_instruction.py | 60 ++++++ 2 files changed, 236 insertions(+) create mode 100644 projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_amax_true16_activation.py diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_amax_true16_activation.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_amax_true16_activation.py new file mode 100644 index 000000000000..ecee36429c84 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_amax_true16_activation.py @@ -0,0 +1,176 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + +"""Regression test for the true16 (real-true16) activation-clamp defect (ROCM-3994). + +hipBLASLt / TensileLite's ``AMaxKernelGenerator`` emits the ``v_max_f16`` +activation clamp for half inputs. Under a compiler that enforces real16 mode +(``+real-true16``) the *fake16* form of this instruction:: + + v_max_f16 v[vgprOutput], v[vgprOutput], abs(v[vgprValue+0]) + +is rejected by the assembler ("operands are not valid for this GPU or mode"). +The correct *true16* form binds a half-word selector (``.l``) to every register +operand, with the suffix placed *inside* the closing paren of ``abs(...)``:: + + v_max_f16 v[vgprOutput].l, v[vgprOutput].l, abs(v[vgprValue+0].l) + +This is a pure host-side codegen assertion: it renders the generator's module to +a string and checks the emitted text. No GPU and no assembler are required. The +target arch is forced from a gfx string, so no real device has to be present. + +RED on ``origin/develop`` (emits fake16), GREEN on the true16 fix branch (emits +true16). The test is parametrized across the whole gfx11 family enabled by the +``NoSDWA`` arch cap (isaVersion 11/12), not a single arch, so it does not repeat +the original mistake of pinning one target. +""" + +import os +import re +import sys + +import pytest + +# Make the tensilelite root importable (mirrors gpu_test_helpers). +SCRIPT_DIR = os.path.dirname(os.path.abspath(__file__)) +TENSILE_ROOT = os.path.abspath(os.path.join(SCRIPT_DIR, "..", "..", "..")) +if TENSILE_ROOT not in sys.path: + sys.path.insert(0, TENSILE_ROOT) + +from gpu_test_helpers import init_rocisa # noqa: E402 + +from Tensile.Common.Architectures import gfxToIsa # noqa: E402 +from Tensile.Common.DataType import DataType # noqa: E402 + +import AMaxGenerator # noqa: E402 + + +# gfx11 family covered by the NoSDWA arch cap (checkInList(isaVersion[0], {11, 12})). +# All of these are wave32 and select the true16 activation path on the fix branch. +GFX11_TARGETS = [ + "gfx1100", + "gfx1101", + "gfx1102", + "gfx1103", + "gfx1150", + "gfx1151", + "gfx1152", + "gfx1153", +] + +# ROCM-3994 known-bug quarantine (W-KNOWN-BUG two-PR flow). +# +# This reproducer lands ahead of the true16 fix (users/ericwan/true16-patch). +# On develop the AMax activation clamp is emitted in the fake16 form, so the +# assertions below fail -> reported as xfailed -> CI stays green. When the fix +# lands, the emitted text becomes true16, the test passes, and strict=True turns +# that XPASS into a hard failure -- forcing the fix PR to delete this marker. +# The test always executes (it is quarantined, not disabled). +# +# Time-box: tracked by ROCM-3994; re-evaluate by the next hipBLASLt release if +# still unfixed. REMOVE this marker in the fix PR. +_ROCM3994_XFAIL = pytest.mark.xfail( + reason="ROCM-3994: fake16 v_max_f16 emitted for the AMax activation clamp; " + "real-true16 assemblers reject it. Remove this marker in the true16 fix PR.", + strict=True, + raises=AssertionError, +) + +# Matches a v_max_f16 whose destination and both sources carry a true16 half-word +# selector (.l/.h). Critically, the abs() source suffix is *inside* the paren. +_TRUE16_RE = re.compile( + r"v_max_f16\s+" + r"v\[[^\]]*\](?:\.[lh]),\s+" # dst.l + r"v\[[^\]]*\](?:\.[lh]),\s+" # src0.l + r"(?:abs\(v\[[^\]]*\](?:\.[lh])\)|v\[[^\]]*\](?:\.[lh]))" # src1(.l) or abs(src1.l) +) + +# The bare fake16 form: a v_max_f16 with no half-word selector on any operand. +_FAKE16_RE = re.compile( + r"v_max_f16\s+" + r"v\[[^\]]*\],\s+" # dst (no .l/.h) + r"v\[[^\]]*\],\s+" # src0 (no .l/.h) + r"(?:abs\(v\[[^\]]*\]\)|v\[[^\]]*\])" # src1 (no .l/.h) +) + + +def _make_generator(target: str) -> "AMaxGenerator.AMaxKernelGenerator": + """Build a half-input AMaxKernelGenerator initialized for ``target``. + + Uses the shared init_rocisa helper (gfxToIsa -> ri.init -> ri.setKernel), + which drives the same rocIsa singleton (``_global_ti``) the generator reads + its arch caps from. gfx11 is wave32. + """ + isa = gfxToIsa(target) + assert isa, f"unknown gfx target: {target}" + init_rocisa(target=target, wavesize=32) + half = DataType("H") + return AMaxGenerator.AMaxKernelGenerator( + i_type=half, + o_type=half, + scale_type=DataType("S"), + num_workitems=256, + num_load_count=4, + num_load_size=4, + wavefront_size=32, + arch=target, + isa=isa, + is_scale=False, + ) + + +def _vmax_lines(text: str): + return [ln.strip() for ln in text.splitlines() if "v_max_f16" in ln] + + +@_ROCM3994_XFAIL +@pytest.mark.unit +@pytest.mark.gfx11 +@pytest.mark.parametrize("target", GFX11_TARGETS) +def test_amax_activation_clamp_is_true16(target): + """max_per_data must emit the true16 v_max_f16 (with abs source).""" + gen = _make_generator(target) + text = str(gen.max_per_data(0)) + + lines = _vmax_lines(text) + assert lines, f"[{target}] expected at least one v_max_f16 in max_per_data:\n{text}" + + for ln in lines: + # Positive: true16 half-word selectors present on the clamp. + assert _TRUE16_RE.search(ln), ( + f"[{target}] v_max_f16 is not in true16 form (missing .l selector " + f"and/or abs suffix inside paren): {ln!r}" + ) + # Negative: reject the bare fake16 form outright. + assert not _FAKE16_RE.search(ln), ( + f"[{target}] v_max_f16 emitted in fake16 form (no half-word " + f"selector) — this is the ROCM-3994 defect: {ln!r}" + ) + + # The abs() source suffix must bind *inside* the paren: abs(v[..].l), never abs(v[..]).l + assert "abs(" in text, f"[{target}] expected an abs() source operand:\n{text}" + assert not re.search(r"abs\(v\[[^\]]*\]\)\.[lh]", text), ( + f"[{target}] true16 suffix placed outside abs() paren (invalid form):\n{text}" + ) + + +@_ROCM3994_XFAIL +@pytest.mark.unit +@pytest.mark.gfx11 +@pytest.mark.parametrize("target", GFX11_TARGETS) +def test_amax_merge_sum_is_true16(target): + """merge_sum must emit the true16 v_max_f16 (plain register sources).""" + gen = _make_generator(target) + text = str(gen.merge_sum()) + + lines = _vmax_lines(text) + assert lines, f"[{target}] expected at least one v_max_f16 in merge_sum:\n{text}" + + for ln in lines: + assert _TRUE16_RE.search(ln), ( + f"[{target}] merge_sum v_max_f16 is not in true16 form: {ln!r}" + ) + assert not _FAKE16_RE.search(ln), ( + f"[{target}] merge_sum v_max_f16 emitted in fake16 form " + f"(ROCM-3994 defect): {ln!r}" + ) diff --git a/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py b/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py index 449e0b18ba0d..e07980049aa2 100644 --- a/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py +++ b/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py @@ -25,6 +25,8 @@ from copy import deepcopy import math +import pytest + def test_instruction_common(): from rocisa.instruction import SMovB32 @@ -430,6 +432,64 @@ def test_instruction_scalar_float(): assert "// mutated" not in str(b) +@pytest.mark.xfail( + reason="ROCM-3994 (W-KNOWN-BUG two-PR flow): the EMaxF16 true16 helper does not " + "exist on develop, so this reproducer fails at import. The true16 fix PR " + "(users/ericwan/true16-patch) must delete this marker; strict=True flips " + "the resulting XPASS to a failure if it is left behind. Time-box: next " + "hipBLASLt release.", + strict=True, + raises=(ImportError, AttributeError), +) +def test_instruction_vmax_f16_true16(): + # Regression for the true16 (real-true16) activation-clamp defect (ROCM-3994). + # + # On a NoSDWA target (isaVersion 11/12) the EMaxF16 helper must emit the + # v_max_f16 activation clamp in true16 form, binding a half-word selector + # (.l) to every register operand. The abs() source keeps its suffix *inside* + # the closing paren: abs(v[..].l) is valid, abs(v[..]).l is not. + # + # RED on develop: EMaxF16 (and the VMaxF16 true16 kwarg) do not exist there, + # so the import / construction below fails outright. GREEN on the fix branch: + # the emitted text carries the .l selectors. + # + # Assertions are substring/regex based to stay agnostic to the ~50-column + # comment padding, matching test_instruction_swait_xcnt / _global_wb. + import re + + import rocisa + from rocisa.container import vgpr + from rocisa.instruction import EMaxF16 + + # Initialize a NoSDWA gfx11 ISA so EMaxF16 selects the true16 path. + isa = (11, 0, 0) + ri = rocisa.rocIsa.getInstance() + ri.init(isa, "amdclang++", False) + ri.setKernel(isa, 32) + assert ri.getArchCaps()["NoSDWA"], "expected NoSDWA cap for gfx11" + + # abs() source (the activation-clamp shape from AMaxGenerator.max_per_data). + inst = EMaxF16(vgpr("Output"), vgpr("Output"), vgpr("Value+0", isAbs=True)) + s = str(inst) + assert re.search( + r"v_max_f16\s+v\[vgprOutput\]\.l,\s+v\[vgprOutput\]\.l,\s+abs\(v\[vgprValue\+0\]\.l\)", + s, + ), f"activation clamp not in true16 form: {s!r}" + # The suffix must be inside the abs() paren, never outside it. + assert "abs(v[vgprValue+0].l)" in s + assert "abs(v[vgprValue+0]).l" not in s + # And it must not be the bare fake16 form. + assert "abs(v[vgprValue+0])," not in s and not s.rstrip().endswith("abs(v[vgprValue+0])") + + # Plain register sources (the merge_sum shape): every operand gets .l. + inst2 = EMaxF16(vgpr("Output"), vgpr("Output"), vgpr("OutputB")) + s2 = str(inst2) + assert re.search( + r"v_max_f16\s+v\[vgprOutput\]\.l,\s+v\[vgprOutput\]\.l,\s+v\[vgprOutputB\]\.l", + s2, + ), f"merge clamp not in true16 form: {s2!r}" + + if __name__ == "__main__": test_instruction_common() test_instruction_cvt() From 02d2875e3aa197869aa31313215778acb14219c4 Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Wed, 22 Jul 2026 13:30:18 -0400 Subject: [PATCH 2/4] test(tensilelite): resolve amdclang++ via ROCM_PATH in true16 rocisa 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 --- .../tensilelite/rocisa/test/test_instruction.py | 14 ++++++++++++-- 1 file changed, 12 insertions(+), 2 deletions(-) diff --git a/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py b/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py index e07980049aa2..f8b6e1285f60 100644 --- a/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py +++ b/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py @@ -455,16 +455,26 @@ def test_instruction_vmax_f16_true16(): # # Assertions are substring/regex based to stay agnostic to the ~50-column # comment padding, matching test_instruction_swait_xcnt / _global_wb. + import os import re + import shutil import rocisa from rocisa.container import vgpr from rocisa.instruction import EMaxF16 - # Initialize a NoSDWA gfx11 ISA so EMaxF16 selects the true16 path. + # Initialize a NoSDWA gfx11 ISA so EMaxF16 selects the true16 path. Resolve + # the assembler via ROCM_PATH first (matching test_mubuf.py::_isa_context) so + # this works when ROCm's bin dir is not on PATH (common in CI / Windows). isa = (11, 0, 0) + rocm_path = os.environ.get("ROCM_PATH", "/opt/rocm") + search_path = os.pathsep.join([ + os.path.join(rocm_path, "bin"), + os.path.join(rocm_path, "lib", "llvm", "bin"), + ]) + assembler = shutil.which("amdclang++", path=search_path) or "amdclang++" ri = rocisa.rocIsa.getInstance() - ri.init(isa, "amdclang++", False) + ri.init(isa, assembler, False) ri.setKernel(isa, 32) assert ri.getArchCaps()["NoSDWA"], "expected NoSDWA cap for gfx11" From 01f85654558f23fe3e9d0254eb15d3a55f4b2b68 Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Wed, 22 Jul 2026 14:32:46 -0400 Subject: [PATCH 3/4] test(tensilelite): document true16 rocisa test mechanic; track toolchain 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 --- .../tensilelite/rocisa/test/test_instruction.py | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py b/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py index f8b6e1285f60..81cc681ee351 100644 --- a/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py +++ b/projects/hipblaslt/tensilelite/rocisa/test/test_instruction.py @@ -478,6 +478,15 @@ def test_instruction_vmax_f16_true16(): ri.setKernel(isa, 32) assert ri.getArchCaps()["NoSDWA"], "expected NoSDWA cap for gfx11" + # How this test works: rocisa instruction objects render their assembly text + # through __str__, so we construct the instruction under test, stringify it, + # and assert on the emitted form. vgpr(name, ...) builds a VGPR operand + # container (symbolic names like "Output"/"Value+0" mirror what the AMax + # generator passes; isAbs=True wraps it in abs()). EMaxF16 is the true16-aware + # max helper under test. Nothing is assembled or run on a device. + # TODO(#9720): decouple this from the toolchain/env init above; a + # string-rendering test should not need a resolved assembler path. + # abs() source (the activation-clamp shape from AMaxGenerator.max_per_data). inst = EMaxF16(vgpr("Output"), vgpr("Output"), vgpr("Value+0", isAbs=True)) s = str(inst) From afd51ff35582454b7320a69d927f08dcaf1c1dd7 Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Fri, 24 Jul 2026 10:36:52 -0500 Subject: [PATCH 4/4] build(hipblaslt): install AMaxGenerator.py into tensilelite test tree 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 --- projects/hipblaslt/CMakeLists.txt | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/projects/hipblaslt/CMakeLists.txt b/projects/hipblaslt/CMakeLists.txt index 98c83f05fc6a..806325af22a8 100644 --- a/projects/hipblaslt/CMakeLists.txt +++ b/projects/hipblaslt/CMakeLists.txt @@ -643,6 +643,15 @@ if(HIPBLASLT_INSTALL_TENSILELITE_TEST_ARTIFACTS AND NOT WIN32) DESTINATION "${_tensilelite_install_dest}/scripts" COMPONENT tests ) + + # Imported directly by unit tests (test_amax_true16_activation.py), so it must + # ship at the tensilelite root, not under scripts/. + install( + FILES + "${_tensilelite_src}/AMaxGenerator.py" + DESTINATION "${_tensilelite_install_dest}" + COMPONENT tests + ) # Preserve executable permissions on entry points install( PROGRAMS