fix(kpack): route gfx1250v0 stepping directory to the gfx1250 bundle

shared/kpack/python/rocm_kpack/database_handlers.py

@@ -121,6 +121,47 @@ def should_move(self, path: Path) -> bool:
     r"/library/(" + _GFX_ARCH_PATTERN.pattern + r")/"
 )
 
+# Known, explicit stepping-directory aliases, scoped to the library subtree
+# they're documented for (keyed by _library_dir). hipBLASLt and hipSPARSELt ship
+# gfx1250 A0 kernels under library/gfx1250v0/, not a separate target -- it
+# belongs to its base arch, matching the arch in the filenames. This mirrors
+# TensileLite's own explicit build-alias table (tensilelite/Common/
+# Architectures.py: ARCH_BUILD_ALIASES), which defines this exact one entry and
+# no others. Deliberately scoped per library and a literal lookup, not a
+# general "any arch + vN" pattern or a cross-library default: only this one
+# case, for these two libraries, is verified. Neither a different stepping,
+# nor a different library (e.g. rocBLAS, which does not produce this
+# directory today), is guaranteed to want the same treatment.
+LIBRARY_ARCH_DIR_STEPPING_ALIASES = {
+    "hipblaslt/library": {"gfx1250v0": "gfx1250"},
+    "hipsparselt/library": {"gfx1250v0": "gfx1250"},
+}
+
+
+def resolve_stepping_alias(
+    preceding_parts: tuple[str, ...], part: str
+) -> Optional[str]:
+    """Return the base arch if `part` is a known stepping-directory alias
+    immediately inside the library subtree it's documented for, else None.
+
+    `preceding_parts` is the path components before `part` (e.g.
+    `("stage", "lib", "hipblaslt", "library")`); a match requires it to end
+    with that library's own path components (e.g. `("hipblaslt", "library")`),
+    not just contain them anywhere.
+
+    Unlike _TensileHandler.detect() (which is scoped to one handler's own
+    _library_dir), this checks every library in
+    LIBRARY_ARCH_DIR_STEPPING_ALIASES: callers like the architecture-
+    separation verifier examine files across every library at once, not
+    through a single handler instance.
+    """
+    for library_dir, aliases in LIBRARY_ARCH_DIR_STEPPING_ALIASES.items():
+        base = aliases.get(part)
+        prefix = tuple(library_dir.split("/"))
+        if base is not None and preceding_parts[-len(prefix) :] == prefix:
+            return base
+    return None
+
 
 class _TensileHandler(DatabaseHandler):
     """Base class for Tensile-based library handlers (rocBLAS, hipBLASLt, hipSPARSELt).
@@ -163,6 +204,12 @@ def detect(self, path: Path, prefix_root: Path) -> Optional[str]:
         if dir_match:
             return dir_match.group(1)
 
+        for alias, base in LIBRARY_ARCH_DIR_STEPPING_ALIASES.get(
+            self._library_dir, {}
+        ).items():
+            if f"/library/{alias}/" in path_str:
+                return base
+
         return None
 
 

shared/kpack/python/rocm_kpack/tools/verify_artifacts.py

@@ -26,6 +26,7 @@
     AotritonHandler,
     ComposableKernelHandler,
     MIOpenHandler,
+    resolve_stepping_alias,
 )
 from rocm_kpack.coff.kpack_transform import (
     HIPF_MAGIC as COFF_HIPF_MAGIC,
@@ -335,6 +336,12 @@ def _check_architecture_separation(self, artifacts: list[Path]) -> None:
                         directory_arch = aotriton.detect(file, artifact)
                         if directory_arch is not None:
                             directory_arches.append(directory_arch)
+                    elif (
+                        alias_base := resolve_stepping_alias(
+                            directory_parts[:index], part
+                        )
+                    ) is not None:
+                        directory_arches.append(alias_base)
                     else:
                         directory_arches.extend(arch_pattern.findall(part))
                 # Handlers may preserve xnack features while filename/directory

shared/kpack/tests/test_database_handlers.py

@@ -193,6 +193,19 @@ def test_detect_arch_subdir_xnack(self, handler, prefix_root):
         result = handler.detect(file_path, prefix_root)
         assert result == "gfx942-xnack+"
 
+    def test_detect_arch_subdir_stepping_not_applicable(self, handler, prefix_root):
+        """The gfx1250v0 stepping alias is scoped to hipBLASLt/hipSPARSELt and
+        must not apply to rocBLAS, even though rocBLAS shares the same
+        _TensileHandler base class. rocBLAS does not produce this directory
+        today, but if it ever did, a directory-only file should stay
+        unclassified rather than being silently merged into gfx1250."""
+        file_path = prefix_root / "lib/rocblas/library/gfx1250v0/TensileManifest.txt"
+        file_path.parent.mkdir(parents=True)
+        file_path.touch()
+
+        result = handler.detect(file_path, prefix_root)
+        assert result is None
+
     def test_reject_wrong_directory(self, handler, prefix_root):
         """Test that files not in rocblas/library are rejected."""
         file_path = prefix_root / "lib/other/library/TensileLibrary_gfx1100.dat"
@@ -335,6 +348,50 @@ def test_detect_arch_subdir_manifest(self, handler, prefix_root):
         result = handler.detect(file_path, prefix_root)
         assert result == "gfx1100"
 
+    @pytest.mark.parametrize(
+        "filename",
+        [
+            "TensileLibrary_lazy_gfx1250.dat.zlib",
+            "TensileLibrary_Type_HH_gfx1250.dat.zlib",
+            "TensileLibrary_Type_HH_gfx1250.co",
+            "Kernels.so-000-gfx1250.hsaco",
+        ],
+    )
+    def test_detect_arch_subdir_stepping(self, handler, prefix_root, filename):
+        """Stepping subdirectory library/gfx1250v0/ belongs to the gfx1250 bundle."""
+        file_path = prefix_root / "lib/hipblaslt/library/gfx1250v0" / filename
+        file_path.parent.mkdir(parents=True)
+        file_path.touch()
+
+        result = handler.detect(file_path, prefix_root)
+        assert result == "gfx1250"
+
+    @pytest.mark.parametrize(
+        "directory",
+        [
+            # Wrong revision: only v0 is a known, verified stepping.
+            "gfx1250v1",
+            # Wrong base arch: only gfx1250 is known to have a v0 stepping.
+            "gfx90av0",
+            "gfx942v0",
+        ],
+    )
+    def test_detect_arch_subdir_unknown_stepping_not_merged(
+        self, handler, prefix_root, directory
+    ):
+        """Only gfx1250v0 is a known, verified stepping alias. Neither a
+        different revision nor a different base arch should be silently folded
+        into a base arch -- a directory-only file (no arch in its own filename)
+        under such a directory should stay unclassified rather than guessed."""
+        file_path = (
+            prefix_root / "lib/hipblaslt/library" / directory / "TensileManifest.txt"
+        )
+        file_path.parent.mkdir(parents=True)
+        file_path.touch()
+
+        result = handler.detect(file_path, prefix_root)
+        assert result is None
+
     def test_reject_wrong_directory(self, handler, prefix_root):
         """Test that files not in hipblaslt/library are rejected."""
         file_path = prefix_root / "lib/rocblas/library/TensileLibrary_gfx1100.dat"
@@ -424,6 +481,18 @@ def test_detect_arch_subdir_dat(self, handler, prefix_root):
         result = handler.detect(file_path, prefix_root)
         assert result == "gfx950"
 
+    def test_detect_arch_subdir_stepping(self, handler, prefix_root):
+        """Stepping subdirectory library/gfx1250v0/ belongs to the gfx1250 bundle."""
+        file_path = (
+            prefix_root
+            / "lib/hipsparselt/library/gfx1250v0/TensileLibrary_lazy_gfx1250.dat.zlib"
+        )
+        file_path.parent.mkdir(parents=True)
+        file_path.touch()
+
+        result = handler.detect(file_path, prefix_root)
+        assert result == "gfx1250"
+
     def test_reject_wrong_directory(self, handler, prefix_root):
         file_path = prefix_root / "lib/rocblas/library/something_gfx942.co"
         file_path.parent.mkdir(parents=True)

shared/kpack/tests/test_strict_identity.py

@@ -78,6 +78,54 @@ def test_split_preserves_only_selected_database_payloads(tmp_path, selected):
     assert not list((dest / "miopen_lib_generic").rglob("*.db.txt"))
 
 
+def test_split_routes_stepping_directory_to_base_arch(tmp_path):
+    # hipBLASLt installs gfx1250 A0 kernels under library/gfx1250v0/ with
+    # gfx1250 filenames; every file there must ship with the gfx1250 bundle.
+    source = tmp_path / "blas_lib_gfx125X-all"
+    prefix = "math-libs/BLAS/hipBLASLt/stage"
+    library = source / prefix / "lib/hipblaslt/library"
+    layout = {
+        "gfx1250": "gfx1250",
+        "gfx1250v0": "gfx1250",
+        "gfx1250-strict": "gfx1250-strict",
+    }
+    for directory, target in layout.items():
+        (library / directory).mkdir(parents=True)
+        for name in [
+            f"TensileLibrary_lazy_{target}.dat.zlib",
+            f"TensileLibrary_Type_HH_{target}.dat.zlib",
+            f"TensileLibrary_Type_HH_{target}.co",
+        ]:
+            (library / directory / name).write_text(directory)
+    write_artifact_manifest(source, [prefix])
+    dest = tmp_path / "split"
+    splitter = ArtifactSplitter(
+        artifact_prefix="blas_lib",
+        toolchain=Toolchain(),
+        database_handlers=[HipBLASLtHandler()],
+        gpu_targets=["gfx1250", "gfx1250-strict"],
+    )
+    splitter.split(source, dest)
+
+    def library_dirs(artifact):
+        output = dest / artifact / prefix / "lib/hipblaslt/library"
+        return {d.name: sorted(p.name for p in d.iterdir()) for d in output.iterdir()}
+
+    assert set(library_dirs("blas_lib_gfx1250")) == {"gfx1250", "gfx1250v0"}
+    assert len(library_dirs("blas_lib_gfx1250")["gfx1250v0"]) == 3
+    assert set(library_dirs("blas_lib_gfx1250-strict")) == {"gfx1250-strict"}
+    assert not [
+        p
+        for p in (dest / "blas_lib_generic").rglob("*")
+        if p.is_file() and "hipblaslt/library" in p.as_posix()
+    ]
+    verifier = ArtifactVerifier(dest, Toolchain())
+    verifier._check_architecture_separation(
+        [dest / "blas_lib_gfx1250", dest / "blas_lib_gfx1250-strict"]
+    )
+    assert verifier.results[-1].passed
+
+
 def test_archive_and_bundler_identity(tmp_path):
     target = "gfx1250-strict"
     assert (
@@ -134,6 +182,55 @@ def test_verifier_detects_target_directory_contamination(tmp_path):
     assert not verifier.results[-1].passed
 
 
+@pytest.mark.parametrize(
+    "artifact_target,passed", [("gfx1250", True), ("gfx1250-strict", False)]
+)
+def test_verifier_maps_stepping_directory_to_base_arch(
+    tmp_path, artifact_target, passed
+):
+    artifact = tmp_path / f"blas_lib_{artifact_target}"
+    directory = artifact / "stage/lib/hipblaslt/library/gfx1250v0"
+    directory.mkdir(parents=True)
+    (directory / "TensileLibrary_lazy_gfx1250.dat.zlib").write_text("data")
+    verifier = ArtifactVerifier(tmp_path, Toolchain())
+    verifier._check_architecture_separation([artifact])
+    assert verifier.results[-1].passed == passed
+
+
+def test_verifier_stepping_alias_scoped_to_known_libraries(tmp_path):
+    """The gfx1250v0 alias is scoped to hipBLASLt/hipSPARSELt and must not
+    apply under rocblas/library/, even though the directory name matches."""
+    artifact = tmp_path / "blas_lib_gfx1250"
+    directory = artifact / "stage/lib/rocblas/library/gfx1250v0"
+    directory.mkdir(parents=True)
+    (directory / "TensileManifest.txt").write_text("data")
+    verifier = ArtifactVerifier(tmp_path, Toolchain())
+    verifier._check_architecture_separation([artifact])
+    assert not verifier.results[-1].passed
+
+
+@pytest.mark.parametrize(
+    "directory",
+    [
+        # Wrong revision: only v0 is a known, verified stepping.
+        "gfx1250v1",
+        # Wrong base arch: only gfx1250 is known to have a v0 stepping.
+        "gfx90av0",
+    ],
+)
+def test_verifier_flags_unknown_stepping_directory(tmp_path, directory):
+    """Only gfx1250v0 is a known, verified exception. A hypothetical, unverified
+    stepping suffix must still be flagged as contamination rather than silently
+    passed through, since we have no evidence it belongs with its base arch."""
+    artifact = tmp_path / "blas_lib_gfx1250"
+    stepping_dir = artifact / "stage/lib/hipblaslt/library" / directory
+    stepping_dir.mkdir(parents=True)
+    (stepping_dir / "TensileManifest.txt").write_text("data")
+    verifier = ArtifactVerifier(tmp_path, Toolchain())
+    verifier._check_architecture_separation([artifact])
+    assert not verifier.results[-1].passed
+
+
 @pytest.mark.parametrize("valid", [True, False])
 def test_verifier_checks_archives_under_component_stage(tmp_path, valid):
     target = "gfx1250-strict"