feat(hip-kernel-provider): add gfx942 dense attention AOT ingestor engine

dnn-providers/hip-kernel-provider/CMakeLists.txt

@@ -483,6 +483,7 @@ if(HIPKERNELPROVIDER_ENABLE_TESTS)
     set(HIPKERNELPROVIDER_ASM_SDPA_INTEGRATION_CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/ASM_SDPA_ENGINE_test_categories_integration.yaml")
     set(HIPKERNELPROVIDER_GFX950_ATTENTION_INTEGRATION_CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/hipkernel_Gfx950Attention_test_categories_integration.yaml")
     set(HIPKERNELPROVIDER_GFX942_ATTENTION_UNIFIED_3D_INTEGRATION_CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/hipkernel_Gfx942AttentionUnified3d_test_categories_integration.yaml")
+    set(HIPKERNELPROVIDER_GFX942_ATTENTION_DENSE_INTEGRATION_CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/hipkernel_Gfx942AttentionDense_test_categories_integration.yaml")
     set(ROCKE_ENGINE_CTEST_CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/ROCKE_ENGINE_test_categories_external.yaml")
     set(HKP_PACK_CTEST_CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/HKP_PACK_test_categories_external.yaml")
     set(HIPKERNELPROVIDER_CTEST_INSTALL_FILE "")
@@ -493,6 +494,7 @@ if(HIPKERNELPROVIDER_ENABLE_TESTS)
        AND EXISTS "${HIPKERNELPROVIDER_ASM_SDPA_INTEGRATION_CATEGORIES_YAML}"
        AND EXISTS "${HIPKERNELPROVIDER_GFX950_ATTENTION_INTEGRATION_CATEGORIES_YAML}"
        AND EXISTS "${HIPKERNELPROVIDER_GFX942_ATTENTION_UNIFIED_3D_INTEGRATION_CATEGORIES_YAML}"
+       AND EXISTS "${HIPKERNELPROVIDER_GFX942_ATTENTION_DENSE_INTEGRATION_CATEGORIES_YAML}"
        AND EXISTS "${ROCKE_ENGINE_CTEST_CATEGORIES_YAML}"
        AND EXISTS "${ROCM_LIBRARIES_ROOT}/shared/ctest/TestCategories.cmake")
         include("${ROCM_LIBRARIES_ROOT}/shared/ctest/TestCategories.cmake")
@@ -503,6 +505,7 @@ if(HIPKERNELPROVIDER_ENABLE_TESTS)
             "${HIPKERNELPROVIDER_ASM_SDPA_INTEGRATION_CATEGORIES_YAML}"
             "${HIPKERNELPROVIDER_GFX950_ATTENTION_INTEGRATION_CATEGORIES_YAML}"
             "${HIPKERNELPROVIDER_GFX942_ATTENTION_UNIFIED_3D_INTEGRATION_CATEGORIES_YAML}"
+            "${HIPKERNELPROVIDER_GFX942_ATTENTION_DENSE_INTEGRATION_CATEGORIES_YAML}"
             "${ROCKE_ENGINE_CTEST_CATEGORIES_YAML}"
             "${HIPKERNELPROVIDER_CTEST_CATEGORIES_YAML}"
         )

dnn-providers/hip-kernel-provider/HIP_MLOPS_ENGINE_test_categories.yaml

@@ -13,18 +13,17 @@
 # See https://github.com/ROCm/rocm-libraries/issues/8852 for background on the
 # category execution model.
 # Tier behaviour (cumulative — each tier is a superset of the ones below it):
-#   - quick: run non-Full tests only.
+#   - quick: run tests with neither a Comprehensive* nor a Full* prefix.
 #   - standard: quick coverage (quick tests). No distinct Standard* prefix
-#     exists, so standard runs the same non-Full tests as quick.
-#   - comprehensive: standard + quick coverage. No distinct Comprehensive*
-#     prefix exists, so comprehensive runs the same non-Full tests as standard.
+#     exists, so standard runs the same tests as quick.
+#   - comprehensive: standard + quick coverage, plus Comprehensive*
+#     parameterised expansions.
 #   - full: all tests (non-Full + Full* parameterised expansions), excluding
 #     only disabled tests.
 #
-# Neither the unit tests in tests/ nor the ASM SDPA / rocKE-client integration
-# tests in integration_tests/ currently use a Full*/Standard*/Comprehensive*
-# GTest prefix, so all four tiers below run the same coverage until a suite
-# adopts one of those prefixes (matching miopen-provider's convention).
+# Comprehensive* and Full* are GTest instantiation prefixes (matching
+# miopen-provider's convention): the gfx942 dense attention knob suite's
+# per-kernel rows are Comprehensive/, and FlyDSL's every-object cases Full/.
 #
 # These categories apply to hip_kernel_provider_tests and
 # hip_kernel_provider_integration_tests. The external hipdnn_integration_tests
@@ -33,31 +32,33 @@
 # GTest naming conventions.
 test_categories:
   quick:
-    description: "Fast sanity checks (target: < 5 min) - excludes Full* parameterised expansions"
+    description: "Fast sanity checks (target: < 5 min) - excludes Comprehensive* and Full* parameterised expansions"
     test_patterns:
       - "*"
     exclude:
       - "*DISABLED*"
+      - "Comprehensive*"
       - "Full*"
     labels:
       - "quick"
       - "pre-commit"
       - "smoke"
 
   standard:
-    description: "Pre-checkin validation (cumulative: runs quick coverage) - excludes Full* parameterised expansions"
+    description: "Pre-checkin validation (cumulative: runs quick coverage) - excludes Comprehensive* and Full* parameterised expansions"
     test_patterns:
       - "*"
     exclude:
       - "*DISABLED*"
+      - "Comprehensive*"
       - "Full*"
     labels:
       - "standard"
       - "pr"
       - "precheckin"
 
   comprehensive:
-    description: "Nightly comprehensive validation (cumulative: runs standard + quick coverage) - excludes Full* parameterised expansions"
+    description: "Nightly comprehensive validation (cumulative: runs standard + quick coverage, plus Comprehensive*) - excludes Full* parameterised expansions"
     test_patterns:
       - "*"
     exclude:
@@ -84,6 +85,7 @@ test_categories:
       - "*"
     exclude:
       - "*DISABLED*"
+      - "Comprehensive*"
       - "Full*"
     labels:
       - "ffm-quick"

dnn-providers/hip-kernel-provider/config/hipkernel_Gfx942AttentionDense.toml

@@ -0,0 +1,17 @@
+# Integration test configuration for hipkernel:Gfx942AttentionDense
+# Filters use GTest-style globs: * matches any sequence of characters
+
+[meta]
+version = 1
+
+# All bundles for this engine are forward-only dense SDPA on gfx942.
+# No tolerance overrides needed: the forward pass uses the common SdpaAttributes
+# tolerance (tol::sdpa::getToleranceFwd<T>()). That includes the FP32 stats (LSE) output
+# of the dense BSHD stats bundles, graded at getToleranceFwd<float>() = 1e-5; the knob
+# suite measures the LSE against the CPU reference at under 4e-6 absolute.
+#
+# No [[test_skips]] entry: the engine's packs arch-prune before the matcher, so a
+# non-gfx942 device declines through the ordinary path. An entry cannot express
+# "skip everywhere except gfx942" -- the schema has no negation, and an empty
+# `archs` list means EVERY arch (TestSettings.hpp: `entry.archs.empty() || ...`),
+# so a skip written that way would silence gfx942 too.

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

@@ -100,11 +100,12 @@ too. Two answers decide which engine-pinned tests are registered:
 
 - **Per-arch registrations** (the census, and the external and gpu_ref integration
   checks) install into a gfx-specific shard and are gated on
-  `hkp_gfx950_attention_dense_available()`, which reads what the shard for that arch
-  ships.
-- **Host-side sources that pin an engine**, such as the gfx950 knobs suite, are gated on
-  `hkp_product_offers_engine()`, which reads what the product root carries for any arch
-  under this build's filters. It never reads `GPU_TARGETS`: the host libraries and
+  `hkp_product_ships_engine_for_arch()`, which reads what the shard for that arch
+  ships. The single-arch rocKE attention engines pinned against the SdpaFwd bundles
+  are registered from one table, `HKP_PINNED_ATTENTION_ENGINES` in `src/CMakeLists.txt`.
+- **Host-side sources that pin an engine**, such as the attention knob suites, are gated
+  on `hkp_product_offers_engine()`, which reads what the product root carries for any
+  arch under this build's filters. It never reads `GPU_TARGETS`: the host libraries and
   tests must work with another build's arch content.
 
 A probe that cannot answer leaves the root wired, so the packer reports what is wrong

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

@@ -1407,24 +1407,6 @@ function(hkp_product_ships_engine_for_arch out_var engine arch)
     endif()
 endfunction()
 
-# ---------------------------------------------------------------------------
-# hkp_gfx950_attention_dense_available(<out>)
-#   hkp_product_ships_engine_for_arch() for hipkernel:Gfx950AttentionDense on gfx950.
-# ---------------------------------------------------------------------------
-function(hkp_gfx950_attention_dense_available out_var)
-    hkp_product_ships_engine_for_arch(_available "hipkernel:Gfx950AttentionDense" gfx950)
-    set(${out_var} "${_available}" PARENT_SCOPE)
-endfunction()
-
-# ---------------------------------------------------------------------------
-# hkp_gfx942_attention_unified_3d_available(<out>)
-#   hkp_product_ships_engine_for_arch() for hipkernel:Gfx942AttentionUnified3d on gfx942.
-# ---------------------------------------------------------------------------
-function(hkp_gfx942_attention_unified_3d_available out_var)
-    hkp_product_ships_engine_for_arch(_available "hipkernel:Gfx942AttentionUnified3d" gfx942)
-    set(${out_var} "${_available}" PARENT_SCOPE)
-endfunction()
-
 # ---------------------------------------------------------------------------
 # hkp_product_offers_engine(<out_var> <engine>)
 #   TRUE when the `product` root, under this build's excluded folders and disabled
@@ -1433,7 +1415,7 @@ endfunction()
 #   Arch-agnostic: it reads what the root offers, never GPU_TARGETS, because it gates
 #   host and test content outside the arch content, which must work with another build's
 #   arch content. Registrations that ship in a per-arch shard use
-#   hkp_gfx950_attention_dense_available() instead.
+#   hkp_product_ships_engine_for_arch() instead.
 #
 #   FALSE when `product` is neither wired nor recorded dormant, and when it was never
 #   probed (an empty root). A probe that could not answer reads as
@@ -1465,16 +1447,6 @@ 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

dnn-providers/hip-kernel-provider/descriptor-packaging/tests/test_attention_bundle_filters.py

@@ -0,0 +1,180 @@
+"""The attention engines' bundle selection agrees between CMake and the category YAMLs.
+
+Each HKP_PINNED_ATTENTION_ENGINES row in src/CMakeLists.txt names the GTEST_FILTER
+variable its engine's external integration check (and GPU-reference target in
+src/integration_tests/) passes, and the ../CMakeLists.txt variable holding its
+hipkernel_*_test_categories_integration.yaml, from which the check's CTest suites take
+theirs under YAML categorization, one per category. Both must name the same bundles, so
+editing one without the other fails here. Rows sharing a YAML must share a filter, and
+every attention categories YAML must be some row's. Each attention knob suite in
+src/integration_tests/CMakeLists.txt pins engines the table names.
+
+A pattern is "<tier>_<bundle>", <tier> one of the four bundle tier directories or "*"
+for all of them. Each category selects every bundle the CMake filter does at each tier
+it runs, and the categories together run every tier the CMake filter does.
+Excludes are not compared: the CMake filters have none, and the YAMLs' (DISABLED,
+SdpaBackward) reach no SdpaFwd bundle.
+"""
+
+import re
+from pathlib import Path
+
+import pytest
+import yaml
+
+pytestmark = pytest.mark.quick
+
+PROVIDER = Path(__file__).resolve().parents[2]
+PROVIDER_CMAKE = PROVIDER / "CMakeLists.txt"
+SRC_CMAKE = PROVIDER / "src/CMakeLists.txt"
+INTEGRATION_CMAKE = PROVIDER / "src/integration_tests/CMakeLists.txt"
+KNOB_SUITES_DIR = PROVIDER / "src/integration_tests/kernel_ingestor_engine"
+PACKS = PROVIDER / "src/tests/engines/kernel_ingestor_engine/packs"
+TABLE = "HKP_PINNED_ATTENTION_ENGINES"
+ROW_FIELDS = 6
+
+
+def pinned_engines():
+    """The table's rows: arch, slug, engine, filter, TOML, categories YAML variable."""
+    text = SRC_CMAKE.read_text(encoding="utf-8")
+    calls = re.findall(rf"\bset\(\s*{TABLE}\s([^)]*)\)", text)
+    assert len(calls) == 1, f"{SRC_CMAKE}: {TABLE} set {len(calls)} times"
+    fields = re.sub(r"#[^\n]*", "", calls[0]).split()
+    assert (
+        fields and len(fields) % ROW_FIELDS == 0
+    ), f"{SRC_CMAKE}: {TABLE} has {len(fields)} fields, not rows of {ROW_FIELDS}"
+    return [fields[i : i + ROW_FIELDS] for i in range(0, len(fields), ROW_FIELDS)]
+
+
+def categories_yaml(var):
+    """The file name ../CMakeLists.txt sets `var` to."""
+    text = PROVIDER_CMAKE.read_text(encoding="utf-8")
+    names = re.findall(
+        rf'\bset\(\s*{var}\s+"\$\{{CMAKE_CURRENT_SOURCE_DIR\}}/([^"/]+)"\s*\)', text
+    )
+    assert len(names) == 1, f"{PROVIDER_CMAKE}: {var} set {len(names)} times"
+    return names[0]
+
+
+def yaml_filters():
+    """{categories YAML: {filter variables of the rows naming it}}."""
+    pairs = {}
+    for *_, filter_var, _toml, yaml_var in pinned_engines():
+        pairs.setdefault(categories_yaml(yaml_var), set()).add(filter_var)
+    return pairs
+
+
+PINNED = yaml_filters()
+
+TIERS = ("quick", "standard", "comprehensive", "full")
+TIERED = re.compile(r"(?P<tier>quick|standard|comprehensive|full|\*)_(?P<bundle>.+)")
+
+
+def cmake_filter(name):
+    """`name`'s value as src/CMakeLists.txt defines it, ${...} references expanded."""
+    text = SRC_CMAKE.read_text(encoding="utf-8")
+    # set(<var> "<arg>"...) or string(JOIN "<glue>" <var> "<arg>"...).
+    calls = re.findall(
+        rf'\b(?:set|string)\(\s*(?:JOIN\s+"([^"]*)"\s+)?{name}\s([^)]*)\)', text
+    )
+    assert len(calls) == 1, f"{SRC_CMAKE}: {name} defined {len(calls)} times"
+    glue, args = calls[0]
+    # A set() list reaches --gtest_filter ':'-joined, as string(JOIN ":") does.
+    value = (glue or ":").join(re.findall(r'"([^"]*)"', args))
+    return re.sub(r"\$\{(\w+)\}", lambda ref: cmake_filter(ref[1]), value)
+
+
+def selection(patterns, where):
+    """{(tier, bundle)} the positive gtest `patterns` select; a "*" tier is all four."""
+    pairs = set()
+    for pattern in patterns:
+        match = TIERED.fullmatch(pattern)
+        assert match, f"{where}: {pattern!r} has no tier prefix ({', '.join(TIERS)}, *)"
+        tiers = TIERS if match["tier"] == "*" else (match["tier"],)
+        pairs.update((tier, match["bundle"]) for tier in tiers)
+    return pairs
+
+
+@pytest.mark.parametrize("row", pinned_engines(), ids=lambda row: f"{row[0]}_{row[1]}")
+def test_each_row_names_an_arch_its_toml_and_its_packs_suite(row):
+    """A misspelt arch or engine local name ships nothing for the row, so its census,
+    external integration check and GPU-reference target drop out silently; fail here."""
+    arch, _slug, local, _filter_var, toml, _yaml_var = row
+    assert re.fullmatch(r"gfx[0-9a-f]+", arch), f"{SRC_CMAKE}: {TABLE} arch {arch!r}"
+    config = PROVIDER / "config" / toml
+    assert config.is_file(), f"{SRC_CMAKE}: {TABLE} {local}: no {config}"
+    suite = PACKS / f"Test{local}Packs.cpp"
+    assert suite.is_file(), f"{SRC_CMAKE}: {TABLE} {local}: no {suite}"
+    # The census keys its expected inventory by arch, so a wrong arch pins none.
+    assert f'{{"{arch}",' in suite.read_text(
+        encoding="utf-8"
+    ), f"{SRC_CMAKE}: {TABLE} {local}: {suite.name} pins no {arch} inventory"
+
+
+def knob_suites():
+    """The _knob_suites rows of src/integration_tests/CMakeLists.txt: (source, engines)."""
+    text = INTEGRATION_CMAKE.read_text(encoding="utf-8")
+    calls = re.findall(r"\bset\(\s*_knob_suites\s([^)]*)\)", text)
+    assert len(calls) == 1, f"{INTEGRATION_CMAKE}: _knob_suites set {len(calls)} times"
+    fields = re.sub(r"#[^\n]*", "", calls[0]).split()
+    assert (
+        fields and len(fields) % 2 == 0
+    ), f"{INTEGRATION_CMAKE}: _knob_suites {fields}"
+    return [(fields[i], fields[i + 1].split(",")) for i in range(0, len(fields), 2)]
+
+
+@pytest.mark.parametrize(
+    "source, engines",
+    [pytest.param(source, engines, id=source) for source, engines in knob_suites()],
+)
+def test_each_knob_suite_pins_engines_the_table_names(source, engines):
+    """A misspelt engine is offered by no product, so its suite compiles silently out."""
+    assert (KNOB_SUITES_DIR / source).is_file(), f"{INTEGRATION_CMAKE}: no {source}"
+    pinned = {local for _arch, _slug, local, *_ in pinned_engines()}
+    unknown = sorted(set(engines) - pinned)
+    assert not unknown, f"{INTEGRATION_CMAKE}: {source} pins {unknown}, not in {TABLE}"
+
+
+def test_rows_sharing_a_categories_yaml_share_a_filter():
+    """One YAML selects one set of bundles, so it can agree with one filter only."""
+    shared = {
+        name: sorted(filters) for name, filters in PINNED.items() if len(filters) > 1
+    }
+    assert not shared, f"{SRC_CMAKE}: {TABLE} rows pair a YAML with filters {shared}"
+
+
+def test_every_attention_category_yaml_is_compared():
+    """A new attention engine's YAML is checked against its filter, not skipped."""
+    found = {
+        path.name
+        for path in PROVIDER.glob("hipkernel_*Attention*_test_categories_*.yaml")
+    }
+    assert found == set(PINNED)
+
+
+@pytest.mark.parametrize(
+    "yaml_name, name",
+    sorted((yaml_name, name) for yaml_name, names in PINNED.items() for name in names),
+)
+def test_categories_select_the_cmake_filter_bundles(yaml_name, name):
+    positive = cmake_filter(name).partition("-")[0]
+    expected = selection(positive.split(":"), f"{SRC_CMAKE} {name}")
+
+    path = PROVIDER / yaml_name
+    categories = yaml.safe_load(path.read_text(encoding="utf-8"))["test_categories"]
+    assert categories, f"{path}: no test_categories"
+    union = set()
+    for category, info in categories.items():
+        selected = selection(info["test_patterns"], f"{path} {category}")
+        tiers = {tier for tier, _ in selected}
+        want = {(tier, bundle) for tier in tiers for _, bundle in expected}
+        assert selected == want, (
+            f"{path} {category}: at tiers {sorted(tiers)}, missing "
+            f"{sorted(want - selected)} and extra {sorted(selected - want)} against "
+            f"{name} in {SRC_CMAKE}"
+        )
+        union |= selected
+    assert union == expected, (
+        f"{path}: categories select {sorted(union)}; "
+        f"{name} in {SRC_CMAKE} selects {sorted(expected)}"
+    )

dnn-providers/hip-kernel-provider/descriptor-packaging/tests/test_hkp_pack_wiring.py

@@ -406,7 +406,7 @@ def answer(shipped, probe_ok):
 set_property(GLOBAL PROPERTY HKP_PACK_PROBE_OK_product {probe_ok})
 set_property(GLOBAL PROPERTY HKP_PACK_SHIPPED_ENGINES_product
     "{_cmake_string(json.dumps(shipped))}")
-hkp_gfx950_attention_dense_available(_available)
+hkp_product_ships_engine_for_arch(_available "hipkernel:Gfx950AttentionDense" gfx950)
 if(_available)
     set(_available TRUE)
 else()
@@ -496,7 +496,7 @@ def gates(force_probe_failure):
             + _wire_root("product", arches="gfx90a", enable_rocke=enable_rocke)
             + f"""{forced}
 hkp_product_offers_engine(_offered "test_fixture:attention")
-hkp_gfx950_attention_dense_available(_available)
+hkp_product_ships_engine_for_arch(_available "hipkernel:Gfx950AttentionDense" gfx950)
 get_property(_dormant GLOBAL PROPERTY HKP_PACK_DORMANT_LABELS)
 {_report("dormant", "${_dormant}")}
 foreach(_name IN ITEMS offered available)