---
name: rocm-hipdnn-superbuild-test
description: "Run tests against an existing hipDNN superbuild. Supports per-component selection (hipdnn, miopen-provider, hipblaslt-provider, hip-kernel-provider, integration-tests), unit/integration/external-integration scope, and gtest filtering. Reproduces the cross-provider external-integration-check suite. Handles Windows DLL PATH automatically. This skill executes targets; for what the cross-provider suite actually runs and how to read its result, defer to the hipdnn-integration-testing skill."
metadata:
  decimalai-skill-id: "b16f0b2f-7bba-4097-8dd6-b7d386b0cfb0"
  decimalai-version-id: "5b707af2-0d6f-42fc-8fa5-a6386c81419f"
  decimalai-version: "7"
  decimalai-body-sha256: "b6e155a30e836c134a1e30b886f4d20b16f2bbbab3e854f812c5e5df64b62a16"
  decimalai-name: "rocm/hipdnn-superbuild-test"
  decimalai-format: "1"
  decimalai-invocation: "model"
  decimalai-category: "testing"
  decimalai-source: "https://app.decimal.ai/s/rocm-hipdnn-superbuild-test@7/SKILL.md"
---
# hipDNN Superbuild Test Runner

Use this skill when the user asks to test an existing hipDNN superbuild. It does not configure or build the project. If no superbuild exists, tell the user to build first with `hipdnn-superbuild`.

## Inputs

Infer options from the user request:

- **Component**: `hipdnn`, `miopen`, `hipblaslt`, `hip-kernel`, `integration-tests`, or `all`; default `all`
- **Scope**: `unit`, `integration`, `external-integration`, or `all`; default `unit`. `external-integration` covers the cross-provider `hipdnn_integration_tests` suite (the `<provider>-external-integration-check` targets)
- **Filter**: optional gtest filter; when present, run test binaries directly
- **Verbose**: use verbose test targets when requested
- **Keep going**: continue after failures only when requested
- **ROCm path**: optional `ROCM_PATH=<path>` override; Linux defaults to `/opt/rocm`
- **Jobs**: optional explicit parallelism only when the user requests it and active workspace instructions permit it

## Workflow

1. Determine the repository root:
   ```bash
   git rev-parse --show-toplevel
   ```

2. Resolve paths:
   - Build directory: honor active workspace instructions first; otherwise use `<repo-root>/build`.
   - Binary directory: `<build-dir>/bin`. `<binary-path>` below is the full path to a test
     executable under it.
   - `<installed-ctest-root>`: the CTest root of an *installed* tree, which for this
     provider is `<install-prefix>/bin/hip_kernel_provider`, not the prefix itself.
   - `<build-ctest-root>`: `<build-dir>/dnn-providers/hip-kernel-provider/src`, the
     build-tree registration of the categorized engine-pinned suites, with absolute paths.
     List one engine's with `ctest --test-dir <build-ctest-root> -N -L <engine-label>`,
     where the label is the engine ID with every character outside `[A-Za-z0-9_.-]` mapped
     to `_` (for example `hipkernel_Gfx950AttentionDense`). The provider's
     `install_CTestTestfile.cmake` (or `CTestTestfile.cmake.install`) in the build tree is
     install-only: CTest does not read it there. The per-arch census shard runs in either
     tree, at
     `<build-dir>/lib/hipdnn_plugins/engines/test_arch_content/hip-kernel-provider/census/<arch>`
     in the build tree (`<build-dir>/bin/...` on Windows) and the same relative path under
     the install prefix.
   - `<PY>`: for the corpus-sweep command below, the Python interpreter the active
     workspace or repository instructions mandate, otherwise the active venv's
     `python`. Resolve it once and substitute the full path.
   - `<GEN>`: the IngestorGenerator root, `<repo-root>/projects/hipdnn/tools/IngestorGenerator`.
   - Helper scripts: skills are host-level, not tied to a repo checkout — **default to the scripts bundled with the skill you were invoked from** (`<skill-directory>/scripts`), even when working inside a repo or worktree. Do NOT run the `<repo-root>/projects/hipdnn/tools/ai/skills/hipdnn-superbuild-test/scripts` copy just because a checkout is present: it can be a stale stub (on `develop`) or an unmerged in-progress version (on a feature branch). Use the source-checkout copy only when actively developing this skill itself to exercise your in-progress edits, when the invoked skill has no bundled `scripts/` directory, or for ingestor engine create/extend work, where the helper revision must match the checkout being built.

3. Verify the superbuild exists:
   ```bash
   ls <build-dir>/build.ninja
   ```
   Stop if it is missing.

4. Resolve ROCm path on Windows:
   ```bash
   python3 <scripts>/windows_rocm_setup.py --repo-root <repo-root> [--rocm-path <path>]
   ```
   Parse `ROCM_PATH=...` from stdout and set `ROCM_BIN=<rocm-path>/bin`. Skip this step on Linux unless the user supplied an override. On Windows, always pass the resolved `ROCM_BIN` to `cmake_run.py` (steps 6-8) via `--rocm-bin`: it is required both for the runtime PATH and for staging the wheel's System32-shadowed DLLs app-local (see Notes).

5. Discover CMake test targets:
   ```bash
   python3 <scripts>/discover_test_targets.py --build-dir <build-dir> --component <component> --scope <scope>
   ```
   The helper prints `<component>:<target>` lines. It also handles the hip-kernel-provider path-qualified target naming. With `--scope external-integration` (or `all`) it emits **one `<component>:<target>` line and one `<component>:command:<cmdline>` line for each registered external suite**. A provider can register several suites. For example, hip-kernel-provider registers one suite for each engine: `hip-kernel-provider-external-integration-check` for `HIP_MLOPS_ENGINE` and `hip-kernel-provider-asm-sdpa-external-integration-check` for `ASM_SDPA_ENGINE`. Do not assume a single line. Each command is the resolved cross-provider `hipdnn_integration_tests` invocation (with `--test-article`/`--test-engine`/`--test-config`) read from the generated `CTestTestfile.cmake`, with any baked-in `--gtest_filter` removed so that you can supply your own. The `--test-engine` value shows which suite a line belongs to.
   If the helper reports that Ninja target discovery failed, treat that as an invalid or stale build directory and stop with the helper's diagnostic. If discovery succeeds but no targets match, report that the requested component or scope is not present in the existing superbuild.

   For an ingestor engine the discovery component is **`hip-kernel`**, not
   `hip-kernel-provider`. A helper's first provider-prefixed command need not be the
   requested engine's registration; inspect the actual installed CTest entry before
   executing it. Replace `<your-bundle-ctest-target>` with the name your own
   registration creates. The gfx942 dense names here are illustrative, not targets to
   copy: unless your build registers
   `hip_kernel_provider_gfx942_attention_dense_gpu_ref_integration_tests`, copying it
   verbatim fails the second command under `--no-tests=error`:
   ```bash
   ctest --test-dir <installed-ctest-root> -N -V \
     -R '^<your-bundle-ctest-target>$'
   ctest --test-dir <installed-ctest-root> --no-tests=error -V \
     -R '^<your-bundle-ctest-target>$'
   ```
   Require your bundle's exact registration, its engine pin
   (`hipkernel:Gfx942AttentionDense` in the illustration), current installed
   executable/plugin/config paths and intended quick/standard cases. Missing
   registration, wrong pin, zero selected cases, all-skipped support or failed numerical
   comparisons fail this gate. A broad component PASS is not exact-engine evidence.

6. Run tests through `cmake_run.py` when no gtest filter is requested:
   ```bash
   python3 <scripts>/cmake_run.py --build-dir <build-dir> --target <target> [--rocm-path <path>] [--rocm-bin <path>] > <log> 2>&1
   ```
   Add `--jobs <N>` only when explicit jobs are both requested and permitted. For verbose mode, append `-verbose` to the target name.

7. Run direct binaries when a gtest filter is requested:
   ```bash
   python3 <scripts>/cmake_run.py --build-dir <build-dir> --binary <binary-path> --gtest-filter "<filter>" [--extra-arg=<flag> ...] [-- <passthrough args>] [--rocm-path <path>] [--rocm-bin <path>] > <log> 2>&1
   ```
   Use the component-to-binary mapping below to choose binaries. `cmake_run.py` accepts arbitrary passthrough flags for the binary: simple values via repeatable `--extra-arg` (use `--extra-arg=--flag` for flag-like values), or an entire flag list after a literal `--`. Passing multiple tokens inside `--binary` is rejected with a clear error.

8. Reproduce the cross-provider external-integration suite (`--scope external-integration`):
   - To run the whole suite exactly as CI does, build the custom target:
     ```bash
     python3 <scripts>/cmake_run.py --build-dir <build-dir> --target <provider>-external-integration-check [--rocm-path <path>] [--rocm-bin <path>] > <log> 2>&1
     ```
   - To run with a custom gtest filter, take the `<component>:command:<cmdline>` line from step 5, run the first token as `--binary` and the rest after `--`, adding your own `--gtest-filter`:
     ```bash
     python3 <scripts>/cmake_run.py --build-dir <build-dir> --binary <hipdnn_integration_tests> -- <--test-article ... --test-engine ... --test-config ...> --gtest_filter=<filter> > <log> 2>&1
     ```
   - **Before reporting the result, load the `hipdnn-integration-testing`
     skill.** This suite is the only place where exit code 0 is often
     meaningless. That skill applies the suite's own documentation to the
     output. This includes the signals that the suite prints but that never
     fail a run.
   - An ingestor engine's explicit-selection C++ cases run inside the shared `hip_kernel_provider_integration_tests` binary (CTest entry `hip_kernel_provider_integration_tests_quick_suite`, which runs every engine's cases); the ingestor RUNBOOK stage 5 gives the filtered command.

9. For every command, keep full output in a log and show only a short tail on failure. Track pass/fail per component. Stop at the first failure unless keep-going was requested. For `external-integration` runs, report the `Passed:`/`Skipped:`/`Failed:` counts from the binary's "TEST COVERAGE SUMMARY". Never report only the exit code. A run that skips 100% of tests is green, but it does not show that the engine still works.

## Direct Binary Mapping

| Component | Unit Binaries | Integration Binaries | External Integration (cross-provider) |
|-----------|---------------|----------------------|----------------------------------------|
| `hipdnn` | `hipdnn_backend_tests`, `hipdnn_frontend_tests`, `hipdnn_data_sdk_tests`, `hipdnn_flatbuffers_sdk_tests`, `hipdnn_plugin_sdk_tests`, `hipdnn_test_sdk_tests` | `hipdnn_public_backend_tests`, `hipdnn_public_frontend_tests`, `hipdnn_backend_logging_shutdown_tests` | — |
| `miopen` | `miopen_plugin_tests` | `miopen_plugin_integration_tests` | `miopen-provider-external-integration-check` (`hipdnn_integration_tests` + `miopen_plugin`, engine `MIOPEN_ENGINE`) |
| `hipblaslt` | `hipblaslt_plugin_tests` | `hipblaslt_plugin_integration_tests` | `hipblaslt-provider-external-integration-check` (`hipdnn_integration_tests` + `hipblaslt_plugin`, engine `HIPBLASLT_ENGINE`) |
| `hip-kernel` | `hip_kernel_provider_tests` | `hip_kernel_provider_integration_tests` | `hip-kernel-provider-external-integration-check` when present |
| `integration-tests` | `hipdnn_integration_tests_unit_tests` | `hipdnn_integration_tests`, `hipdnn_gpu_ref_tests` | — |

The exact article/engine/config for the external suite is resolved at build time. Use
discovery for available commands, then inspect the specific CTest registration when
proving a named engine; a generic `command:` line is not that proof.

## Packaging probes

When the build was configured with `HIPKERNELPROVIDER_ENABLE_PACKAGING_PROBES=ON`
(superbuild-only, OFF by default), the probe ctest entries are `hkp-probe-<arch>`, one per
architecture explicitly named under the production descriptor root that `GPU_TARGETS`
does not include (automatic; none when the build targets them all), `hkp-probe-<NAME>`,
one per named `hkp_add_packaging_probe()` line in `probes.cmake`, and `hkp-probe-tools`. A full ctest run includes them;
to run them alone, select with the anchored `-R '^hkp-probe-'` and `--no-tests=error`.
`-DHIPKERNELPROVIDER_MIN_PACKAGING_PROBES=N` fails configure when fewer than N probes are
declared (default 0).

In the superbuild CI lanes the probes run once, inside the "Run tests" ctest, with N=1.
A pass does not cover device behavior,
install staging, UKDs the probe does not pack, or standalone-UKD references.

## Ingestor proof boundaries

[The ingestor RUNBOOK](../hipdnn-ingestor-engine/RUNBOOK.md) is the sole ordered
create/extend workflow and states how each gate is invoked. What matters here is what a
pass does not establish: `device_probe.py` success, in either mode, is not dispatch.

Native host proof executes actual typed provider registrations and descriptor loading,
then checks the finalized emitted inventory. Use a fresh process, explicit
`HIPDNN_TEST_EXPECTED_ARCH` from configured packaging architectures, the corresponding
shard and a nonempty exact host-test selection. Missing or unknown architecture
selection, wrong-arch data, absent/extra identities and wrong runtime source kind fail;
packaged runtime source kind is KPACK. Source-text symbol matching and structural
validation cannot certify native hooks, and host loading cannot prove device dispatch.

Numerical acceptance needs a capable independent reference for the actual graph, and
neither SDPA reference supports a sink UID. Record **BLOCKED** when no capable reference
exists; a skip, automatic fallback exhaustion or unverified golden output cannot pass.

The corpus sweep interface is `<PY> <GEN>/tools/sweep.py --config <absolute-YAML>` with
`configs/reduction.sweep.yaml.example`. Correctness is separate from timing,
engine attribution is exact, and resume is bound to current input content. `SWEEP_DONE`
is validated completion; explicit `correctness.enabled: false` yields
`SWEEP_TIMING_ONLY`, never final acceptance; unmet gates yield `SWEEP_INCOMPLETE`. After
tuning or regeneration, repeat artifact/native/device/corpus gates against the final
installation and complete the per-corpus runtime outcome join. Passing only unchanged
baseline cases cannot establish that an extension's new variant served.

## Report

Summarize per-component results:

```text
hipdnn:
  hipdnn-unit-check: PASS
miopen-provider:
  miopen-provider-unit-check: FAIL (see <log>)
```

If a requested component has no matching target, say that it was not present in the existing superbuild and name the preset or component likely needed.

## Notes

- `scripts/cmake_run.py`, `scripts/discover_test_targets.py`, `scripts/windows_rocm_setup.py`, and `scripts/stage_shadowed_dlls.py` are bundled in this skill so linked and copied installs work independently.
- Windows DLL loading is handled by `cmake_run.py`, which sets PATH in Python's subprocess environment before launching CMake or test binaries.
- Windows shadowed-DLL staging: before launching any target or binary on Windows, `cmake_run.py` stages the wheel's `amd_comgr.dll` and HIP runtime `amdhip64_<N>.dll` into `<build-dir>/bin` (via `stage_shadowed_dlls.py`) so the process does not load the driver's stale `System32` copies. This happens on every Windows run, not just for a specific kernel path. Stale comgr breaks MIOpen's runtime JIT (GCN-assembly Winograd solvers are the common failure: `[BuildAsm] comgr status = ERROR` / `unknown emulation: no-xnack`); a stale HIP runtime makes the wheel's rocBLAS fault with `SEH exception with code 0xc0000005` in MIOpen's GEMM conv solvers (e.g. `GemmFwd1x1_0_1`), after which the run can hang. This needs `--rocm-bin` to be passed. A copy is skipped when the staged DLL already matches the wheel's (PE version, else content hash), so it adds no cost on repeat runs. Disable with `--no-stage-shadowed-dlls` if ever needed. To confirm which comgr loaded, run a test with `MIOPEN_LOG_LEVEL=7 MIOPEN_ENABLE_LOGGING=1` and grep for `COMgr v.` (a low version indicates the stale System32 copy; the wheel's is newer). The driver's System32 HIP runtime announces itself with a `HIP Library Path: C:\windows\SYSTEM32\amdhip64_<N>.dll` line at startup; the wheel's runtime prints no such line. MIOpen's `AmdRocmMetadataVersionDetect` HIP version is the build-time version and does not tell you which runtime loaded.
- The build's `stage_shadowed_rocm_dlls` target (`projects/hipdnn/cmake/WindowsDllStaging.cmake` and `dnn-providers/cmake/WindowsDllStaging.cmake`) is the primary mechanism for app-local staging. `cmake_run.py`'s staging is kept on purpose rather than as a leftover: it covers build trees configured before that target existed, and a newly discovered System32-shadowed DLL can be added to `stage_shadowed_dlls.py` right away, ahead of the matching CMake change. When you add a DLL to one, add it to the other.
- Integration tests require an AMD GPU. Unit scope is the default for CPU-only validation.
- Prefer running test binaries through `cmake_run.py` (it wires PATH/ROCM_PATH for the loader); pass extra binary flags via `--extra-arg`/`-- <args>` rather than folding them into `--binary`.

## See also

- `hipdnn-integration-testing` skill — the doc-driven reference behind the
  `external-integration` scope. It loads the suite's documentation
  (`dnn-providers/integration-tests/docs/`) and applies it to what a run did
  and printed. **Load it before interpreting any result from step 8, and
  before concluding an engine did or did not regress.** This skill executes;
  that one explains.