[nunnikri]

Motivation

On A0-revision gfx1250 silicon, hipBLASLt/hipSPARSELt's loader reads exclusively from
library/gfx1250v0/ — by design, with no fallback to library/gfx1250/, since base
gfx1250 kernels give wrong results on A0. kpack's artifact splitter was silently
dropping the Tensile index from that directory during packaging, so every GEMM
dispatch on A0 failed at load with HIPBLAS_STATUS_INVALID_VALUE. This is the
stopgap fix to keep plain gfx1250 builds working on A0 during the ongoing
gfx1250-strict transition, per the explicit requirement that plain gfx1250
builds must not regress in the meantime.

Technical Details

• library/gfx1250v0/ ships a complete Tensile index
  (TensileLibrary_lazy_gfx1250.dat.zlib, etc.) alongside its .co/.hsaco kernel
  files, but kpack's splitter never assigned the index to any architecture bundle:
  the per-arch directory regex in database_handlers.py didn't recognize the v0
  stepping suffix, and the index files end in .dat.zlib, which the filename-based
  check also doesn't match. The .co/.hsaco files survived (filename-matched);
  the index didn't, so nothing could select/load them.
• LIBRARY_ARCH_DIR_STEPPING_ALIASES is an explicit lookup table — keyed by
  library ("hipblaslt/library", "hipsparselt/library") and, within each, by the
  exact stepping directory name ("gfx1250v0") — mapping to the base arch
  ("gfx1250") it belongs to. Deliberately not a general "any arch + vN"
  pattern: hipBLASLt currently defines only gfx1250v0 as a revision subtree
  (mirroring TensileLite's own ARCH_BUILD_ALIASES table, which has this exact one
  entry and no others), and a future stepping directory — a different revision, a
  different base architecture, or a different library — is not guaranteed to want
  the same treatment, so it's intentionally left unmatched (and flagged, not
  silently merged) until verified.
• Scoped per library specifically because _TensileHandler (the base class this
  alias lookup lives on) is shared by rocBLAS as well as hipBLASLt/hipSPARSELt.
  Without the per-library scoping, a hypothetical rocblas/library/gfx1250v0/
  would have been silently merged into the gfx1250 bundle too, even though the
  alias is only documented and verified for hipBLASLt/hipSPARSELt. rocBLAS doesn't
  produce that directory today, so this wasn't a live regression, but it was the
  same category of unverified assumption this fix otherwise removes.
• verify_artifacts.py's architecture-separation checker shares the same lookup
  via a resolve_stepping_alias() helper in database_handlers.py, so the alias
  data and its matching rules live in exactly one place. The verifier's version
  checks that the stepping directory is immediately inside the correct library's
  own path components (a plain tuple-suffix comparison), not just present anywhere
  in the path — this replaced an earlier, less readable version based on raw
  path-component index arithmetic.
• The two directories (gfx1250/ and gfx1250v0/) stay physically separate on
  disk; only which published artifact they ship inside changes. gfx1250-strict
  is unaffected — it remains its own, fully separate artifact.
• Root-cause analysis, the original regex approach, and initial test coverage are
  from George Tseng (hipBLASLt); this PR carries that fix forward into
  rocm-systems, narrowed from a general v<N> regex to the explicit,
  library-scoped alias table above per review feedback.

Issue Tracking

JIRA ID: ROCM-30193

Test Plan

• Ran the kpack test suite (pytest shared/kpack/tests/, excluding
  test_wheel_splitter.py, which fails to collect on an unrelated missing
  rocm_bootstrap dependency pre-existing in this environment).
• HipBLASLtHandler/HipSparseLtHandler tests cover directory-stepping detection,
  including the .dat.zlib index case that was previously dropped.
• Negative tests confirm the alias does not apply, and the directory stays
  unclassified (handler) or is flagged as contamination (verifier), for: a wrong
  revision (gfx1250v1), a wrong base architecture (gfx90av0, gfx942v0), and
  the wrong library (rocblas/library/gfx1250v0/).
• A full split-pipeline test asserts gfx1250v0/ ships inside the gfx1250
  artifact end-to-end (splitter + verifier), and that gfx1250-strict stays
  separate.
• Verified the remaining failing tests are pre-existing and unrelated by
  re-running the full suite against a clean checkout (same files reverted) and
  confirming an identical failure set.

Test Result

614 passed, 32 failed, 5 skipped (full suite including pre-existing/unrelated
failures) — identical failure set with and without this change, confirming no
regressions. All new/updated tests for this fix pass.

[therock-pr-bot]

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[marbre]

Thanks @nunnikri, this looks sound, however, we hit this for the first time with gfx1250v0 and I am not sure a regex is the way to go. Could we eventually make the stepping mapping explicit for gfx1250v0 instead of accepting any v<N> suffix? hipBLASLt currently defines only gfx1250v0 as a revision subtree, and TensileLite models it as an explicit build alias. Its architecture parser also warns that names like gfx1250v1 can look valid to ISA parsing without being supported targets.

The current regex would silently route any matching gfx*v<N> directory to its base architecture. That works for gfx1250v0, but could misroute a future stepping directory whose kernels need a separate bundle. An explicit mapping would match the producer’s current contract and avoid assuming all revision suffixes share the base architecture’s packaging target.

[nunnikri]

Thanks @nunnikri, this looks sound, however, we hit this for the first time with gfx1250v0 and I am not sure a regex is the way to go. Could we eventually make the stepping mapping explicit for gfx1250v0 instead of accepting any v<N> suffix? hipBLASLt currently defines only gfx1250v0 as a revision subtree, and TensileLite models it as an explicit build alias. Its architecture parser also warns that names like gfx1250v1 can look valid to ISA parsing without being supported targets.

The current regex would silently route any matching gfx*v<N> directory to its base architecture. That works for gfx1250v0, but could misroute a future stepping directory whose kernels need a separate bundle. An explicit mapping would match the producer’s current contract and avoid assuming all revision suffixes share the base architecture’s packaging target.

Thanks @marbre for the review. Updated accordingly

[geotseng-amd]

LGTM. Verified end to end:

• Real-install split (local hipBLASLt + hipSPARSELt install, gpu_targets=[gfx1250, gfx1250-strict]):
  • Base b7db38b: hipBLASLt gfx1250v0/ keeps only 218/437 files; all 219 .dat.zlib, including TensileLibrary_lazy_gfx1250.dat.zlib, land in blas_lib_generic. hipSPARSELt keeps 65/131. The verifier flags 283 files as "gfx1250v". This matches the nightly symptom in ROCM-30193.
  • This PR (c9f67f7): all 437 + 131 gfx1250v0 files land in blas_lib_gfx1250; gfx1250/ and gfx1250-strict/ are unchanged; the architecture-separation check passes (1152 / 572 files).
• kpack suite: 503 -> 516 passed, failure set identical to base (pre-existing env-only failures).
• The new tests fail on the unfixed source (5 failures, exactly the bug cases), so they guard the regression.
• The explicit alias matches TensileLite's ARCH_BUILD_ALIASES, which has the single entry gfx1250v0 -> target gfx1250. Agree with marbre that this is better than a generic v regex.

Nit (non-blocking): the PR description's Technical Details still describes the optional v regex from the first commit. Please update it to the explicit alias table before squash-merge.

[marbre]

Thanks for updating the PR accordingly. My review agent pointed out the following:

The explicit alias is applied outside the library layouts that use it. The map is consulted by the shared _TensileHandler (/develop/rocm-systems/shared/kpack/python/rocm_kpack/database_handlers.py:178), which also handles rocBLAS. As a result, a file under rocblas/library/gfx1250v0/ can be assigned to the gfx1250 bundle, although the documented and tested alias is for hipBLASLt and hipSPARSELt. The verifier also normalizes any path component named gfx1250v0, without checking that it is under a relevant library/ directory (verify_artifacts.py (/develop/rocm-systems/shared/kpack/python/rocm_kpack/tools/verify_artifacts.py:339)).

I recommend scoping the alias to the relevant library handlers and directory context, so this exception can’t classify or validate unrelated paths.

As of today rocBLAS doesn't produce that directory thus this shouldn't be an issue today and it would not introduce a current regression. Thus this feedback is not blocking but restricting the alias to hipBLASLt and hipSPARSELt would be a tidy scope improvement. Wdyt?

[nunnikri]

Thanks for updating the PR accordingly. My review agent pointed out the following:

The explicit alias is applied outside the library layouts that use it. The map is consulted by the shared _TensileHandler (/develop/rocm-systems/shared/kpack/python/rocm_kpack/database_handlers.py:178), which also handles rocBLAS. As a result, a file under rocblas/library/gfx1250v0/ can be assigned to the gfx1250 bundle, although the documented and tested alias is for hipBLASLt and hipSPARSELt. The verifier also normalizes any path component named gfx1250v0, without checking that it is under a relevant library/ directory (verify_artifacts.py (/develop/rocm-systems/shared/kpack/python/rocm_kpack/tools/verify_artifacts.py:339)).
I recommend scoping the alias to the relevant library handlers and directory context, so this exception can’t classify or validate unrelated paths.

As of today rocBLAS doesn't produce that directory thus this shouldn't be an issue today and it would not introduce a current regression. Thus this feedback is not blocking but restricting the alias to hipBLASLt and hipSPARSELt would be a tidy scope improvement. Wdyt?

Yes i agree on that. This change should be very specific to the existing v0 folders. Updated the same in latest commit. Thanks again for your time on the review

[marbre]

No major blockers. Code and tests look good to me but I think your agent has been a bit to verbose and we can cut down the comment / docstring a bit.

Please update the PR description / squashed commit message before landing as this is still giving the state when you implemented the regex and it does not reflect the current state.

[systems-assistant]

TheRock Submodule Bump Activity

Newest first

• 2026-10-22 — Included in TheRock via ROCm/TheRock#9425.