Repository navigation
revert(hipblaslt): remove the gfx1250v0 subtree and the asicRevision check - #15123
Open
geotseng-amd wants to merge 1 commit into
Open
geotseng-amd wants to merge 1 commit into
geotseng-amd wants to merge 1 commit into
Conversation
⏳ 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:
|
…check The V0 stepping is now served only by gfx1250-strict: with HSA_DISABLE_GFX12_STRICT=0 it reports gfx1250-strict and loads library/gfx1250-strict/. hipBLASLt and hipSPARSELt no longer read hipDeviceProp_t::asicRevision, and the build no longer produces library/gfx1250v0/ (gfx1250-strict's logic built for the gfx1250 target). This undoes the gfx1250v0 parts of #11777 by hand; later commits touched the same files, so a git revert does not apply. The xnack handling built on top of it (compilerTargetOf, the xnack- capability overrides, outDirs) is kept. - tensile_host.cpp in hipBLASLt and hipSPARSELt, and validate_library_layout.py, are restored to their state before the gfx1250v0 change. - HIPBLASLT_BUILD_GFX1250V0 and TensileCreateLibrary --gfx1250v0 are removed, and with them ARCH_BUILD_ALIASES, tuningArchOf, deviceTargetFeaturesOf, withArchBuildAliases and archRenames. - The layout validator reports a library/gfx1250v0/ left in a reused install prefix. - Comments that named gfx1250v0 as the V0 build now name gfx1250-strict. Behavior changes: - A V0 part at HSA_DISABLE_GFX12_STRICT=1 or unset reports gfx1250 and loads library/gfx1250/ without an error. V0 needs =0 and a package that ships gfx1250-strict/. - A gfx1250-only build no longer runs --check-all on gfx1250-strict logic. Testing (FFM, ROCm 10.3.0a20261103): - gfx1250 and gfx1250-strict code objects match develop except for StinkyTofu's random label hashes; hipSPARSELt is byte-identical. - V0 topology at =0 loads gfx1250-strict/: hipblaslt-test *quick* 18043 passed, hipsparselt-test *quick* 36816 passed, tox Tests/common 198 passed and 0 failed. - TensileLite unit tests and both pre-commit configs pass. Co-authored-by: Cursor <cursoragent@cursor.com>
geotseng-amd
force-pushed
the
users/geotseng/revert-gfx1250v0
branch
from
November 5, 2026 16:45
74ff18e to
7ab98b9
Compare
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #15123 +/- ##
========================================
Coverage 77.75% 77.75%
========================================
Files 2904 2903 -1
Lines 491207 491131 -76
Branches 74462 74447 -15
========================================
- Hits 381923 381869 -54
+ Misses 84288 84266 -22
Partials 24996 24996
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
Contributor
|
Is it guaranteed that all V0 deployments set HSA_DISABLE_GFX12_STRICT=0? Otherwise, V0 may silently load gfx1250 kernels and produce incorrect results. |
jichangjichang
approved these changes
Nov 6, 2026
This branch has not been deployed
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 : AIHPBLAS-4727
JIRA ID : AIHPBLAS-4728
JIRA ID : AIHPBLAS-4729
JIRA ID : AIHPBLAS-4730
JIRA ID : AIHPBLAS-4768
JIRA ID : AIHPSPAR-275
JIRA ID : ROCM-31226
JIRA ID : ROCM-30511
Motivation
hipBLASLt and hipSPARSELt load their GPU kernels from one directory per GPU architecture under
library/. The gfx1250 chip has two hardware revisions (steppings). The early one, V0, reports the namegfx1250-strictwhen the ROCm runtime variableHSA_DISABLE_GFX12_STRICTis0, and its kernels are built intolibrary/gfx1250-strict/.#11777 also kept a second route for V0, added in its commit 6fdee38 ("build v0 kernels for gfx1250 and pick them by ASIC revision"). That route is a third directory,
library/gfx1250v0/, holding the same V0 kernels built for thegfx1250compiler target. Both libraries picked that directory by reading the chip's revision number (hipDeviceProp_t::asicRevision) whenever V0 reported plaingfx1250. That route costs an extra TensileLite build pass (on our host the hipBLASLt build took 538 s with it and 322 s without). It also keeps a revision check in two libraries and an architecture-alias layer in TensileLite. This PR removes the route, so the namegfx1250-strictis the only thing that selects V0 kernels.Technical Details
This PR undoes the gfx1250v0 parts of 6fdee38, which #11777 squash-merged into develop. About 30 later commits edited the same files, so
git revertdoes not apply; the parts are undone by hand. Those later commits built XNACK support on the same helpers. XNACK is a GPU memory-retry mode named in the target string, as ingfx1250:xnack-, and that support is kept.tensile_host.cppin hipBLASLt and hipSPARSELt is restored byte-for-byte to its form before 6fdee38. Neither library readsasicRevisionnow; the device name alone pickslibrary/<name>/.rocblaslt_arch_revision.hppandarch_revision_gtest.cppare deleted.HIPBLASLT_BUILD_GFX1250V0andTensileCreateLibrary --gfx1250v0are removed.-a "gfx1250;gfx1250-strict"now produces two directories instead of three.ARCH_BUILD_ALIASESis removed, along with the helpers that existed only for it:tuningArchOf,deviceTargetFeaturesOf,withArchBuildAliasesand thearchRenamesparameter.Toolchain/Source.pyreturns to its earlier output paths, which give the same result once no alias exists.validate_library_layout.pyis restored to its form before 6fdee38. It now reports alibrary/gfx1250v0/left over in a reused install prefix, and a new test intest_validate_library_layout.pychecks that.test_arch_steppings.pyare deleted. Two tests are added. One checks thatcompilerTargetOfreturns every target unchanged when no compiler is given. The other checks thatTensileCreateLibrarypasses the namegfx1250-strictto the StinkyTofu instruction scheduler, which turns on V0-only workarounds by that name.gfx1250v0the V0 build now saygfx1250-strict.Two behaviors change:
HSA_DISABLE_GFX12_STRICT=1or unset reportsgfx1250. It now loadslibrary/gfx1250/, which holds kernels for the production stepping, and prints no error. Before this PR it loadedlibrary/gfx1250v0/, or stopped with an error if that directory was missing. V0 users must set the variable to0and use a package that containslibrary/gfx1250-strict/. TheRock packages gain that directory in ROCm/TheRock#8350, and this PR should merge after it.gfx1250no longer compiles the gfx1250-strict tuning files, so it no longer checks them. Namegfx1250-strict, or useall, to cover them.Other architectures are not affected.
Test Plan
Three environments. Each started from deleted build trees.
FFM simulator. FFM is a software model of an MI450-class gfx1250 GPU. These runs used ROCm 10.3.0a20261103 and compared a clean build of develop
29e5b0a059ewith a clean build of this branch.-a "gfx1250;gfx1250-strict". Compare every code object's.text,.rodataand note sections, and its symbol list, against develop.HSA_DISABLE_GFX12_STRICTset to0, set to1, and unset (5 combinations in all). For each, record withstracewhichlibrary/subdirectory each library opens, and run 2 small tests per library. Repeat with a leftoverlibrary/gfx1250v0/copied in, and again withlibrary/gfx1250-strict/removed.=0:hipblaslt-testandhipsparselt-testwith--gtest_filter='*quick*'; the TensileLite config tests-m common -k "not streamk".-m 'unit and not gpu'),tools/scripts/tests, the rocisa adaptor'stest_emission_consistency.py, andpre-commitwith the root andprojects/hipblasltconfigs.Real V0 hardware. One A0 GPU (
ASIC Revision: 0), with ROCm 10.3.0a20261105 (the newest nightly on 2026-11-05) in a fresh venv. The run used this branch at74ff18e9f2andHSA_DISABLE_GFX12_STRICT=0. It followed every section ofdocs/gfx1250-strict-a0-install.mdexcept 8.7, which covers the removedgfx1250v0path.rocminfo,amdgpu-arch,rocm_agent_enumeratorandinvoke get-gpu-archall reportgfx1250-strict, and that no/opt/rocmpath is left in the environment.ctest), hipBLASLt and hipSPARSELt with-a "gfx1250;gfx1250-strict". Check the ELF machine flags, and check the layout withvalidate_library_layout.py.lddthat each test binary loads the library just built, not the copy in the ROCm SDK.hipblaslt-test,*quick*, the FP4 bench and the mxf4 smoke test, the fullhipsparselt-test, the tox config tests-m common -k "not streamk"(tox 4.67.4), the TensileLite unit tests,test_arch_steppings.py,tools/scripts/tests, and both pre-commit configs over29e5b0a059e..HEAD.Real V1 hardware. One V1 GPU (
ASIC Revision: 1), with the same ROCm 10.3.0a20261105 in a fresh venv and the same sections ofdocs/gfx1250-strict-a0-install.md. The run used this branch after it was rebased onto develop307f8f6d352, at7ab98b94ce9. Compared with74ff18e9f2, the rebase changed one test docstring, intest_r7_createlib_deep_char.py: develop's new_parseLogicFileForLibrarywrapper had been described as takingarchRenames, which this PR removes. ROCr adds-strictonly at revision 0, so V1 reportsgfx1250and the GPU suites runlibrary/gfx1250/;gfx1250-strict/is built and checked but not run.rocminfo,amdgpu-archandinvoke get-gpu-archreportgfx1250atHSA_DISABLE_GFX12_STRICTset to0, set to1, and unset.ldd, andLD_DEBUG=libson a running test, which also lists libraries opened at run time. Record withstracewhichlibrary/subdirectory hipSPARSELt opens at=0and=1, and run the hipBLASLt smoke at=0,=1and unset.307f8f6d352with the same nightly in a separate worktree, and compare its code objects with this branch's.hipblaslt-testcase alone, 3 times on each.Test Result
FFM simulator
library/holdsgfx1250/(454 files) andgfx1250-strict/(440 files); hipSPARSELt has 131 files in each. Neither.socontains the stringgfx1250v0.=0opensgfx1250-strict/. V0 at=1or unset, and revision 2 in both modes, opengfx1250/. All 10 runs (2 tests each) pass. A leftovergfx1250v0/is never opened, and the install check reports 441 violations for it. Withgfx1250-strict/removed, V0 at=0stops withCannot read .../TensileLibrary_lazy_gfx1250-strict..., the same as develop.hipblaslt-test *quick*: 18043 passed, 0 failed.hipsparselt-test *quick*: 36816 passed, 0 failed.-k "not streamk"and selects nothing, so this run used the same pytest command tox builds, invoked directly.setuptoolsin the venv,HIP_PATHset in the shell, and 5 order-dependent cases intest_subtile_ds_read_emit.py.test_arch_steppings.py: 909 passed.tools/scripts/tests: 16 passed. Adaptor emission tests: 300 passed. Both pre-commit configs pass; the hipBLASLt test hook ran 12537 tests with 0 failures.Real V0 hardware
gfx1250-strict, withASIC Revision: 0(0x0).-j 128and ran 2 TensileLite passes,[gfx1250]and[gfx1250-strict]. Files: 454 and 440. Code objects: 224 and 217, with ELF flags0x49and0xEB. The layout check reports OK, and there is nogfx1250v0/directory. hipSPARSELt took 151 s and has 131 files in each directory. StinkyTofuctest: 1860 of 1860 passed.ldd:libhipblaslt,libtensilelite-hostandlibhipsparseltresolve to the new build's install directories, and HIP and HSA to the 10.3.0a20261105 SDK. Both tests printDevice ID 0 : ... gfx1250-strict. hipBLASLt loadslibrary/gfx1250-strict.hipblaslt-test: 69985 tests ran; 69784 passed, 0 failed, 201 skipped (1216 s).*quick*: 18043 passed, 0 failed.nan, 0infand 0 failed. mxf4 smoke: 40 passed.hipsparselt-test: 165405 passed, 0 failed. The first 137973 tests ran in one process. The other 27432 ran in four shards on the same GPU, selected by name. The names of passing tests match--gtest_list_testsexactly: none missing, none run twice.test_subtile_ds_read_emit.py; that file passes 8 of 8 when run alone.test_arch_steppings.py: 909 passed.tools/scripts/tests: 16 passed. Both pre-commit configs: every hook Passed or Skipped, and no file changed.Real V1 hardware
rocminfo,amdgpu-archandinvoke get-gpu-archreportgfx1250at=0,=1and unset, withASIC Revision: 1(0x1). No/opt/rocmpath orHIP_CLANG_PATHis left in the environment.-j 32and ran 2 TensileLite passes,[gfx1250]and[gfx1250-strict]. Files: 454 and 440. Code objects: 224 for targetgfx1250and 217 forgfx1250-strict. The layout check reports OK; there is nogfx1250v0/directory, andlibhipblaslt.sodoes not contain the string. hipSPARSELt took 127 s and has 131 files in each directory, with nogfx1250v0/. StinkyTofuctest: 1874 of 1874 passed.307f8f6d352: the file names match in both directories, and 872 of 894 files are byte-identical. The other 22 have the same instruction addresses and encodings; only the random suffixes StinkyTofu adds to local labels differ.ldd:libhipblaslt.so.1andlibtensilelite-host.so.1resolve tohipblaslt-install/lib,libhipsparselt.so.0tohipsparselt-install/lib, and every other ROCm library to the 10.3.0a20261105 SDK; none resolves to/opt/rocm. UnderLD_DEBUG=libs, onehipblaslt-testcase initializes 34 shared objects, split the same way. The tox client loadslibtensilelite-host.so.1from its own build directory.Device ID 0 : ... gfx1250and loadslibrary/gfx1250at=0,=1and unset; the smoke filter passes 1152 tests at each. hipSPARSELt opens onlylibrary/gfx1250at=0and=1; 384 tests pass at each.hipblaslt-test: 69848 passed, 2 failed, 206 skipped (1239 s). Both failures arenightly_matmul_fast_check_largef32 cases at16385×16383×4097, batch 2, oneTNand oneTT, with wrong values in part of D. Run alone, both passed 3 times on this branch and 3 times on develop307f8f6d352. During the full run, the TensileLite unit tests were also using the GPU.*quick*: 18043 passed, 0 failed.nan, 0infand 0 failed. mxf4 smoke: 40 passed.hipsparselt-test: 165501 passed, 0 failed (2030 s). It ran as four gtest shards (GTEST_TOTAL_SHARDS=4) on the same GPU at once: 41376, 41375, 41375 and 41375 tests.bf16_CLS_gfx1250.yaml, which develop added in83c00bc1056, and it also fails on develop307f8f6d352. Each run, one CLS solution out of 361 at(4095,4095,1,1024)had wrong values, and it was a different solution each time: 48, 78 and 66 wrong values out of 16769025 in three runs on this branch, and 844 on develop.not gpu) and 14405 passed (gpu), 0 failed.test_arch_steppings.py: 909 passed.tools/scripts/tests: 16 passed.307f8f6d352..HEAD: every hook Passed or Skipped, and no file changed. The hipBLASLt test hook reads only staged files, and nothing was staged, so it selected no tests here; the full unit run above covers them.Not run
HSA_DISABLE_GFX12_STRICT=1or unset, which is the behavior change; this was checked only on the simulator.-a alland single-stepping builds.Automated GitHub checks are pending.
Risk level
4/5. This PR changes which kernel directory a V0 part loads when
HSA_DISABLE_GFX12_STRICTis not0, in two libraries, and that change was checked only on the simulator. On real V0 hardware at=0, the full hipBLASLt and hipSPARSELt suites pass with 0 failures and loadlibrary/gfx1250-strict/. On real V1 hardware, both libraries loadlibrary/gfx1250/at every setting of the variable, and the kernels match develop's. The fullhipsparselt-testpasses with 0 failures. The fullhipblaslt-testhad 2 failures, and both passed when rerun alone on this branch and on develop. For every other configuration, the kernels and the directory each library loads were unchanged in the measurements above.Submission Checklist