feat(hipdnn): rocke conv gfx950 dgrad bf16/fp16 engine and ranker

.github/workflows/hipdnn-superbuild-ci.yml

@@ -253,6 +253,22 @@ jobs:
           set -euo pipefail
           dvc pull --recursive dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors
 
+      # The gfx950 dgrad AOT instance lists are DVC-tracked too (same `ingestor`
+      # remote), and hipdnn_ingestor_generator_pytest reads them straight from
+      # the source tree -- without this pull, only the .dvc pointers are present
+      # and every test that loads them fails with FileNotFoundError.
+      - name: Fetch dgrad instance lists
+        run: |
+          set -euo pipefail
+          dvc pull --recursive projects/hipdnn/tools/IngestorGenerator/configs
+          # A pull can exit 0 having fetched nothing, or only one of the two dtype
+          # lists; check both so a partial pull is not mistaken for a full one.
+          bf16_count=$(find projects/hipdnn/tools/IngestorGenerator/configs -name '*_bf16.instances.json' | wc -l)
+          fp16_count=$(find projects/hipdnn/tools/IngestorGenerator/configs -name '*_fp16.instances.json' | wc -l)
+          test "$bf16_count" -gt 0 || { echo "dvc pull produced no bf16 dgrad instance lists" >&2; exit 1; }
+          test "$fp16_count" -gt 0 || { echo "dvc pull produced no fp16 dgrad instance lists" >&2; exit 1; }
+          echo "Fetched $bf16_count bf16 and $fp16_count fp16 dgrad instance list(s)"
+
       # Must precede configure: configuration fails if pytest is not importable
       # by the interpreter descriptor-packaging registers its tests under, and,
       # with HIPDNN_REQUIRE_UHD_GEN_TESTS=ON, if uhd_gen's requirements are not.
@@ -531,6 +547,21 @@ jobs:
           test "$count" -gt 0 || { echo "dvc pull produced no rocKE KDPs" >&2; exit 1; }
           echo "Fetched $count rocKE KDP(s)"
 
+      # Same reasoning as the Windows lane: the gfx950 dgrad AOT instance lists
+      # are DVC-tracked (same `ingestor` remote), and hipdnn_ingestor_generator_pytest
+      # reads them straight from the source tree.
+      - name: Fetch dgrad instance lists
+        run: |
+          set -euo pipefail
+          LOGNAME=github-runner dvc pull --recursive projects/hipdnn/tools/IngestorGenerator/configs
+          # A pull can exit 0 having fetched nothing, or only one of the two dtype
+          # lists; check both so a partial pull is not mistaken for a full one.
+          bf16_count=$(find projects/hipdnn/tools/IngestorGenerator/configs -name '*_bf16.instances.json' | wc -l)
+          fp16_count=$(find projects/hipdnn/tools/IngestorGenerator/configs -name '*_fp16.instances.json' | wc -l)
+          test "$bf16_count" -gt 0 || { echo "dvc pull produced no bf16 dgrad instance lists" >&2; exit 1; }
+          test "$fp16_count" -gt 0 || { echo "dvc pull produced no fp16 dgrad instance lists" >&2; exit 1; }
+          echo "Fetched $bf16_count bf16 and $fp16_count fp16 dgrad instance list(s)"
+
       # Must precede configure: configuration fails if pytest is not importable
       # by the interpreter descriptor-packaging registers its tests under, and,
       # with HIPDNN_REQUIRE_UHD_GEN_TESTS=ON, if uhd_gen's requirements are not.

.pre-commit-config.yaml

@@ -105,7 +105,13 @@ repos:
           ^dnn-providers/hip-kernel-provider/\.clang-format$|
           ^dnn-providers/integration-tests/\.clang-format$
       - id: check-added-large-files
-        exclude: *opted_out_but_scanned
+        # The opted-out projects above, plus the UHD pipeline's trained model binaries
+        # under a descriptor set's heuristics/. Each is regenerated, never edited by hand.
+        exclude: |
+          (?x)
+          projects/hipblaslt/.*|
+          shared/tensile/.*|
+          ^dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/descriptors/.*/heuristics/.*/model\.bin$
 
   - repo: https://github.com/psf/black
     rev: 25.12.0

dnn-providers/hip-kernel-provider/descriptor-packaging/README.md

@@ -526,9 +526,26 @@ export AMD_COMGR_CACHE_DIR=/tmp/comgr-cache   # RAM disk or local disk
 
 `HKP_PACK_JOBS` is read by a **direct child run** of `hkp_pack`. Inside the build it comes
 from the `PACK_JOBS <n>` argument of each `hkp_wire_pack_target()` call: the test roots
-pass `1`, or `2` for `integration`, and the production root and the packaging probes pass
-none, so the packer uses its default. The reason is in the `hkp_wire_pack_target()` header
-in `cmake/HkpPackaging.cmake`.
+pass `1`, or `2` for `integration`, the packaging probes pass none, and the production root
+passes `HIPKERNELPROVIDER_PRODUCT_PACK_JOBS`. That cache variable is empty by default, which
+passes none, so the packer uses its default; set it to size the production pool for a large
+catalog on a wide host. The reason is in the `hkp_wire_pack_target()` header in
+`cmake/HkpPackaging.cmake`.
+
+## rocKE toolchain pins
+
+Two `CACHE` variables, resolved once in `_hkp_resolve_rocke_args()` under
+`HIPKERNELPROVIDER_ENABLE_ROCKE=ON`:
+
+| Variable | Forwarded as | Reaches | Effect |
+|---|---|---|---|
+| `HIPKERNELPROVIDER_ROCKE_COMGR_LIB` | `ROCKE_COMGR_LIB` | the pack step and the ctest entries | Explicit `libamd_comgr` for the rocKE producer to load. Empty by default (derived from the `amd_comgr` package); set it to override, e.g. on Windows where a System32 `amd_comgr.dll` can shadow the ROCm one. |
+| `HIPKERNELPROVIDER_ROCKE_LLVM_FLAVOR` | `ROCKE_LLVM_FLAVOR` | the pack step only | LLVM IR flavor rocKE lowers to (`llvm20`, `llvm22`, `llvm23`, `llvm24`). Empty by default (derived from the loaded COMGR's LLVM); set it to lower exactly as a catalog was swept -- the gfx950 conv AOT catalogs were measured as `llvm22`. An unrecognized value is only rejected when the pack step runs, not at configure. |
+
+Changing either repacks every rocKE-sourced catalog with a different lowering, so a silent
+loss of the forwarding would repack with whatever rocke derives instead, unnoticed.
+`tests/test_hkp_pack_wiring.py` and `tests/rocke/test_hkp_python_environment.py` pin the
+forwarding.
 
 ## Packaging probes
 

dnn-providers/hip-kernel-provider/descriptor-packaging/cmake/HkpPackaging.cmake

@@ -254,15 +254,20 @@ function(_hkp_pack_producer)
     # ROCKE_CPP_QUIET_FALLBACK is deliberately unset: silencing that warning
     # hides the degradation these pins exist to catch.
     #
-    # ROCKE_COMGR_LIB overrides a shadowed System32 amd_comgr on Windows; forward
-    # it when set (runtime resolution, no find_library).
+    # ROCKE_COMGR_LIB overrides a shadowed System32 amd_comgr on Windows, or names the
+    # COMGR whose LLVM a catalog was measured with; forward it when set (runtime
+    # resolution, no find_library). ROCKE_LLVM_FLAVOR pins the IR flavor rocke lowers
+    # to instead of the one it derives from that COMGR's LLVM.
     set(_tool_env "ROCKE_BACKEND=python" "ROCKE_CPP_STRICT=1")
     if(ARG_PACK_JOBS)
         list(APPEND _tool_env "HKP_PACK_JOBS=${ARG_PACK_JOBS}")
     endif()
     if(ARG_ROCKE_COMGR_LIB)
         list(APPEND _tool_env "ROCKE_COMGR_LIB=${ARG_ROCKE_COMGR_LIB}")
     endif()
+    if(ARG_ROCKE_LLVM_FLAVOR)
+        list(APPEND _tool_env "ROCKE_LLVM_FLAVOR=${ARG_ROCKE_LLVM_FLAVOR}")
+    endif()
 
     set(_interp "${ARG_ROCKE_INTERP}" PARENT_SCOPE)
     set(_interp_what "rocKE wheel interpreter (root '${ARG_NAME}')" PARENT_SCOPE)
@@ -311,7 +316,8 @@ endfunction()
 #               ARCHES <list> HIPCC <path>
 #               ROCM_KPACK_DIR <dir> OUT_ROOT <dir> ENABLE_ROCKE <bool>
 #               [ROCKE_INTERP <path> ROCKE_READY <path> ROCKE_PYTHON_DIR <dir>
-#                ROCKE_WHEEL_STAMP <path> [ROCKE_COMGR_LIB <path>]]
+#                ROCKE_WHEEL_STAMP <path> [ROCKE_COMGR_LIB <path>]
+#                [ROCKE_LLVM_FLAVOR <flavor>]]
 #               [EXCLUDE_FOLDERS <name>...] [PACK_JOBS <n>])
 #   Wire the compile -> prune -> pack DAG for ONE authored source root.
 #
@@ -348,8 +354,9 @@ endfunction()
 #   depends on it rather than only the interpreter, so changing a kernel under
 #   rocke/library restages the pack. ROCKE_WHEEL_STAMP is the wheel content
 #   digest, recorded into each rocKE UKD's provenance so a shipped kernel names
-#   the wheel that produced it. ROCKE_COMGR_LIB, if set, is forwarded to the
-#   tool environment.
+#   the wheel that produced it. ROCKE_COMGR_LIB and ROCKE_LLVM_FLAVOR, if set, are
+#   forwarded to the tool environment; both decide the code a rocKE kernel compiles to,
+#   so they are part of the pack's input set and changing either repacks.
 #
 #   OFF: every ROCKE_* keyword is a configure error, since each would name a
 #   toolchain the build does not have. The root runs under the base
@@ -366,7 +373,8 @@ endfunction()
 #   serial path. Omitted, the packer sizes its pool against the machine. Roots have
 #   no ordering edge between them, so the generator runs them at once: the test
 #   roots name a small cap so their pools do not multiply, and the product root
-#   omits it because it is the root expected to be large enough to repay a full pool.
+#   omits it because it is the root expected to be large enough to repay a full pool,
+#   unless HIPKERNELPROVIDER_PRODUCT_PACK_JOBS names one.
 #
 #   NAME is the source label written into every descriptor's provenance. NAME, the
 #   absolute SOURCE_ROOT, OUT_ROOT and ARCHES go into a global registry read by
@@ -375,7 +383,7 @@ endfunction()
 function(hkp_wire_pack_target)
     set(_one NAME SOURCE_ROOT ARCHES HIPCC ROCM_KPACK_DIR
         OUT_ROOT ENABLE_ROCKE ROCKE_INTERP ROCKE_READY ROCKE_PYTHON_DIR
-        ROCKE_COMGR_LIB ROCKE_WHEEL_STAMP PACK_JOBS)
+        ROCKE_COMGR_LIB ROCKE_LLVM_FLAVOR ROCKE_WHEEL_STAMP PACK_JOBS)
     cmake_parse_arguments(PARSE_ARGV 0 ARG "" "${_one}" "EXCLUDE_FOLDERS")
     if(ARG_UNPARSED_ARGUMENTS)
         message(FATAL_ERROR
@@ -456,6 +464,12 @@ function(hkp_wire_pack_target)
     set(_manifest_paths ${_source_inputs} ${_tool_sources})
     list(SORT _manifest_paths)
     string(REPLACE ";" "\n" _manifest_body "${_manifest_paths}")
+    # The COMGR and the IR flavor decide the code every rocKE kernel compiles to, so a
+    # change of either has to repack like a changed source.
+    if(ARG_ENABLE_ROCKE)
+        string(APPEND _manifest_body
+               "\n# rocke comgr: ${ARG_ROCKE_COMGR_LIB}\n# rocke llvm flavor: ${ARG_ROCKE_LLVM_FLAVOR}")
+    endif()
     # cmake-lint: disable=E1126
     #   cmake-lint carries no form spec for file(CONFIGURE) and reports it as an
     #   invalid discriminator. It is valid CMake from 3.18; the floor here is 3.25.
@@ -1168,6 +1182,32 @@ packaging dormant.")
     endif()
 endfunction()
 
+# ---------------------------------------------------------------------------
+# _hkp_resolve_product_pack_jobs(<out_var>)
+#   Declare HIPKERNELPROVIDER_PRODUCT_PACK_JOBS and resolve it to the arguments the
+#   production root is wired with: `PACK_JOBS <n>` when it is set, nothing when it is
+#   empty, so the packer sizes its own pool. Only the production root reads it; the test
+#   roots keep their pinned caps. Anything but empty or a positive integer is fatal here
+#   rather than at pack time, where HKP_PACK_JOBS would reject it mid-build.
+# ---------------------------------------------------------------------------
+function(_hkp_resolve_product_pack_jobs out_var)
+    set(HIPKERNELPROVIDER_PRODUCT_PACK_JOBS "" CACHE STRING
+        "Worker processes the production descriptor pack may spawn (HKP_PACK_JOBS for \
+the production root only). Empty lets the packer size its pool against the machine.")
+
+    set(${out_var} "" PARENT_SCOPE)
+    if("${HIPKERNELPROVIDER_PRODUCT_PACK_JOBS}" STREQUAL "")
+        return()
+    endif()
+    if(NOT HIPKERNELPROVIDER_PRODUCT_PACK_JOBS MATCHES "^[0-9]+$"
+       OR HIPKERNELPROVIDER_PRODUCT_PACK_JOBS LESS 1)
+        message(FATAL_ERROR
+            "hkp: HIPKERNELPROVIDER_PRODUCT_PACK_JOBS must be empty or an integer of 1 "
+            "or more, got '${HIPKERNELPROVIDER_PRODUCT_PACK_JOBS}'.")
+    endif()
+    set(${out_var} PACK_JOBS "${HIPKERNELPROVIDER_PRODUCT_PACK_JOBS}" PARENT_SCOPE)
+endfunction()
+
 # ---------------------------------------------------------------------------
 # _hkp_disabled_families(<out_var>)
 #   The folders of HKP_DESCRIPTOR_FAMILIES whose option is OFF in this build. An entry
@@ -1425,6 +1465,16 @@ function(hkp_product_offers_engine out_var engine)
     endif()
 endfunction()
 
+# ---------------------------------------------------------------------------
+# hkp_gfx950_conv_dgrad_aot_available(<out>)
+#   The same predicate for the shape-generic AOT-instance backward-data engine,
+#   hipkernel:Gfx950ConvDgradAot, which ships its own bundle (the `product` root).
+# ---------------------------------------------------------------------------
+function(hkp_gfx950_conv_dgrad_aot_available out_var)
+    hkp_product_ships_engine_for_arch(_ships "hipkernel:Gfx950ConvDgradAot" gfx950)
+    set(${out_var} "${_ships}" PARENT_SCOPE)
+endfunction()
+
 # ---------------------------------------------------------------------------
 # _hkp_resolve_rocke_args(out_args out_comgr_lib)
 #   Resolve the rocKE toolchain once and return the keyword list every pack target is
@@ -1444,6 +1494,13 @@ the ROCm one. rocke treats the value as the first CANDIDATE and falls through \
 when it does not load, so configure asserts that the library which loaded is \
 the one named here.")
 
+    set(HIPKERNELPROVIDER_ROCKE_LLVM_FLAVOR "" CACHE STRING
+        "LLVM IR flavor the rocKE producer lowers to (llvm20, llvm22, llvm23, llvm24), \
+forwarded into ROCKE_LLVM_FLAVOR for the pack step. Leave it empty to let rocke derive \
+it from the LLVM of the COMGR it loads. Set it to lower exactly as a catalog was swept: \
+the gfx950 conv AOT catalogs were measured as llvm22 on a COMGR of LLVM 22 or newer. An \
+unrecognized value is only rejected when the pack step actually runs, not at configure.")
+
     # Runs before the copy below and before the probe, so the derived value is
     # what both of them see. An explicitly-set value is left alone.
     hkp_default_rocke_comgr_lib()
@@ -1480,6 +1537,9 @@ the one named here.")
     if(_rocke_comgr_lib)
         list(APPEND _rocke_args ROCKE_COMGR_LIB "${_rocke_comgr_lib}")
     endif()
+    if(HIPKERNELPROVIDER_ROCKE_LLVM_FLAVOR)
+        list(APPEND _rocke_args ROCKE_LLVM_FLAVOR "${HIPKERNELPROVIDER_ROCKE_LLVM_FLAVOR}")
+    endif()
 
     set(${out_args} "${_rocke_args}" PARENT_SCOPE)
     set(${out_comgr_lib} "${_rocke_comgr_lib}" PARENT_SCOPE)
@@ -1723,6 +1783,8 @@ function(hkp_add_packaging)
     _hkp_disabled_families(_exclude_folders)
     list(APPEND _root_args EXCLUDE_FOLDERS ${_exclude_folders})
 
+    _hkp_resolve_product_pack_jobs(_product_jobs_args)
+
     # Production descriptors.
     _hkp_wire_root(
         NAME product
@@ -1731,6 +1793,7 @@ function(hkp_add_packaging)
         HIPCC "${HKP_HIPCC}"
         ROCM_KPACK_DIR "${_rocm_kpack_dir}"
         OUT_ROOT "${HIPKERNELPROVIDER_DESCRIPTOR_BUILD_DIR}"
+        ${_product_jobs_args}
         ${_root_args})
 
     # Test descriptors, one pack per authored set. The shared root is packed into both

dnn-providers/hip-kernel-provider/descriptor-packaging/python/hkp_pack/provenance_sidecar.py

@@ -45,10 +45,14 @@
 _READABLE_MAJOR = 1
 _READABLE_MINOR = 0
 
-# A ceiling on one sidecar's decompressed size, far above any real pack: a sidecar
-# holds a few kilobytes per UKD, and a gzip body can inflate a thousandfold, so an
-# uncapped read lets a small corrupt or hostile file exhaust memory.
+# A ceiling on one sidecar's decompressed size: a gzip body can inflate a
+# thousandfold, so an uncapped read lets a small corrupt or hostile file exhaust
+# memory. A sidecar holds a few kilobytes per UKD, so the ceiling grows with the
+# UKDs of the descriptor being attached -- MAX_DECOMPRESSED_BYTES_PER_UKD each,
+# far above any real entry -- and never drops below MAX_DECOMPRESSED_BYTES. A
+# fixed ceiling would refuse a catalog of tens of thousands of kernels.
 MAX_DECOMPRESSED_BYTES = 64 * 1024 * 1024
+MAX_DECOMPRESSED_BYTES_PER_UKD = 64 * 1024
 
 
 def sidecar_name(descriptor_name: str) -> str:
@@ -170,14 +174,19 @@ def detach(descriptor_name: str, doc: dict) -> tuple[str, bytes]:
     return sidecar_name(descriptor_name), encode(sidecar)
 
 
-def _inflate(data: bytes, path: Path) -> bytes:
-    """One gzip member, refused past MAX_DECOMPRESSED_BYTES or with trailing bytes."""
+def decompressed_limit(ukd_count: int = 0) -> int:
+    """The most bytes a sidecar for a descriptor of `ukd_count` UKDs may inflate to."""
+    return max(MAX_DECOMPRESSED_BYTES, ukd_count * MAX_DECOMPRESSED_BYTES_PER_UKD)
+
+
+def _inflate(data: bytes, path: Path, limit: int) -> bytes:
+    """One gzip member, refused past `limit` bytes or with trailing bytes."""
     inflater = zlib.decompressobj(16 + zlib.MAX_WBITS)
-    raw = inflater.decompress(data, MAX_DECOMPRESSED_BYTES + 1)
-    if len(raw) > MAX_DECOMPRESSED_BYTES:
+    raw = inflater.decompress(data, limit + 1)
+    if len(raw) > limit:
         raise HkpPackError(
-            f"provenance sidecar {path} inflates past "
-            f"{MAX_DECOMPRESSED_BYTES} bytes; no packed sidecar is that large"
+            f"provenance sidecar {path} inflates past {limit} bytes; no packed "
+            "sidecar for its descriptor is that large"
         )
     if not inflater.eof:
         raise HkpPackError(f"provenance sidecar {path} is truncated")
@@ -205,9 +214,12 @@ def _require_version(version, path: Path) -> None:
         )
 
 
-def load_sidecar(path) -> dict:
+def load_sidecar(path, ukd_count: int = 0) -> dict:
     """The validated sidecar document at `path`; an unreadable or malformed file,
-    or one of a version this reader does not read, is an HkpPackError."""
+    or one of a version this reader does not read, is an HkpPackError.
+
+    `ukd_count`, the UKDs of the descriptor the sidecar belongs to, sizes the
+    decompression ceiling (see `decompressed_limit`)."""
     path = Path(path)
     try:
         data = path.read_bytes()
@@ -216,7 +228,7 @@ def load_sidecar(path) -> dict:
     except OSError as exc:
         raise HkpPackError(f"cannot read provenance sidecar {path}: {exc}") from exc
     try:
-        sidecar = json.loads(_inflate(data, path))
+        sidecar = json.loads(_inflate(data, path, decompressed_limit(ukd_count)))
     except HkpPackError:
         raise
     except (zlib.error, ValueError, RecursionError) as exc:
@@ -297,7 +309,7 @@ def attach(
             f"{path}: packed descriptor has no provenance sidecar; expected "
             f"{side}.{hint}"
         )
-    sidecar = load_sidecar(side)
+    sidecar = load_sidecar(side, len(ukds))
     where = str(side)
     if sidecar.get("kdp_id") != kdp_id:
         raise HkpPackError(