Skip to content

feat(flags): add HIPDNN_ENABLE_KERNEL_INGESTOR build flag - #7194

Merged
SamuelReeder merged 2 commits into
mainfrom
users/sareeder/almiopen-2399-kernel-ingestor-flag
Aug 10, 2026
Merged

SamuelReeder merged 2 commits into
mainfrom
users/sareeder/almiopen-2399-kernel-ingestor-flag

Conversation

@SamuelReeder

@SamuelReeder SamuelReeder commented Aug 7, 2026 •

Copy link
Copy Markdown
Contributor

Summary

Adds a single opt-in build flag that gates the generic kernel ingestor build-time logic (dynamic engine loading, kpack bundling/packaging, dynamic engine registration) so that work tracked under the hipDNN generic kernel ingestor epic can land incrementally without affecting default builds. See RFC 0017: Universal Kernel Descriptors (rocm-libraries#9533) for the platform this flag gates.

JIRA ID : ALMIOPEN-2399

Risk Assessment

Minimal risk. This declares an inert, default-OFF build flag in FLAGS.cmake with no accompanying source changes; no consuming code currently reads the propagated variable, so default builds and installed artifacts are byte-for-byte unaffected.

ASIC Coverage

No ASIC impact. The flag is plumbing only (default OFF, no gated logic exists yet in hipDNN/hipkernelprovider to exercise); standard PR CI is sufficient.

Testing Summary

  • CMake configure smoke test confirming the flag registers and reports correctly with its default value.
  • CMake configure smoke test confirming the flag can be toggled on via the command line and propagates through therock_finalize_flags()/therock_report_flags() without error.

Testing Checklist

  • CMake configure (flag default) - cmake -S . -B build -DTHEROCK_ENABLE_ALL=OFF -DTHEROCK_ENABLE_HIPDNN=OFF - Status: Passed (reports HIPDNN_ENABLE_KERNEL_INGESTOR = OFF)
  • CMake configure (flag enabled) - cmake -S . -B build -DTHEROCK_ENABLE_ALL=OFF -DTHEROCK_ENABLE_HIPDNN=OFF -DTHEROCK_FLAG_HIPDNN_ENABLE_KERNEL_INGESTOR=ON - Status: Passed (reports HIPDNN_ENABLE_KERNEL_INGESTOR = ON)
  • PR CI - GitHub PR checks - Status: Pending

Technical Changes

  • Declares THEROCK_FLAG_HIPDNN_ENABLE_KERNEL_INGESTOR (default OFF) in FLAGS.cmake.
  • Scopes CMAKE_VARS HIPDNN_ENABLE_KERNEL_INGESTOR=ON to the hipDNN and hipkernelprovider subprojects, matching the epic's plan to gate plugin-sdk and hip-kernel-provider dynamic-engine logic behind one flag.

Declares THEROCK_FLAG_HIPDNN_ENABLE_KERNEL_INGESTOR (default OFF) in FLAGS.cmake, propagating HIPDNN_ENABLE_KERNEL_INGESTOR=ON to the hipDNN and hipkernelprovider subprojects when enabled. This gates the generic kernel ingestor build-time logic (dynamic engine loading, kpack bundling/packaging) tracked under the hipDNN generic kernel ingestor epic, so that work can land incrementally without affecting default builds.

JIRA ID : ALMIOPEN-2399
@therock-pr-bot

therock-pr-bot Bot commented Aug 7, 2026 •

Copy link
Copy Markdown

✅ All Policy Checks Passed

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass PR does not contain code files — Unit Test auto-passed
🚫 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?

@therock-pr-bot

therock-pr-bot Bot commented Aug 7, 2026

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

@SamuelReeder SamuelReeder self-assigned this Aug 7, 2026
@SamuelReeder
SamuelReeder marked this pull request as ready for review August 7, 2026 16:08
@SamuelReeder
SamuelReeder marked this pull request as draft August 7, 2026 16:09
@SamuelReeder
SamuelReeder marked this pull request as ready for review August 7, 2026 16:23
SamuelReeder added a commit to ROCm/rocm-libraries that referenced this pull request Aug 10, 2026
Declares the option in both projects/hipdnn/CMakeLists.txt and dnn-providers/hip-kernel-provider/CMakeLists.txt (default OFF) so it's a single, pre-reserved receiver for TheRock's forwarded flag (see TheRock FLAGS.cmake / ROCm/TheRock#7194), before parallel follow-on work (ALMIOPEN-2400/2401/2402) lands the actual gated code and risks each re-declaring the option independently.

JIRA ID : ALMIOPEN-2399
@SamuelReeder
SamuelReeder merged commit 69c37f9 into main Aug 10, 2026
149 of 158 checks passed
@SamuelReeder
SamuelReeder deleted the users/sareeder/almiopen-2399-kernel-ingestor-flag branch August 10, 2026 16:13
@github-project-automation github-project-automation Bot moved this from TODO to Done in TheRock Triage Aug 10, 2026
SamuelReeder added a commit to ROCm/rocm-libraries that referenced this pull request Aug 10, 2026
## Summary

Declares `HIPDNN_ENABLE_KERNEL_INGESTOR` (default `OFF`) in both
`projects/hipdnn/CMakeLists.txt` and
`dnn-providers/hip-kernel-provider/CMakeLists.txt`, matching the
variable TheRock's `FLAGS.cmake` forwards to both subprojects
([ROCm/TheRock#7194](ROCm/TheRock#7194)). This
reserves a single option ahead of the parallel follow-on stories
(ALMIOPEN-2400, ALMIOPEN-2401, ALMIOPEN-2402) that will land the actual
`#ifdef`/`if()`-gated kernel-ingestor code, so those PRs don't each
independently declare — and risk conflicting on — the same option. See
[RFC 0017: Universal Kernel
Descriptors](projects/hipdnn/docs/rfcs/0017_UniversalKernelDescriptor.md)
for the platform this gates.

JIRA ID : ALMIOPEN-2399

## Risk Assessment

Minimal risk. Adds an inert, default-OFF `option()` in two CMakeLists
files with no gated code yet; no existing build path, target, or test is
affected.

## ASIC Coverage

No ASIC impact. Declaration-only change with no compiled/runtime
behavior difference at any flag value.

## Testing Summary

- Standalone CMake configure of `projects/hipdnn` and
`dnn-providers/hip-kernel-provider` with
`-DHIPDNN_ENABLE_KERNEL_INGESTOR=ON`, confirming the option registers in
`CMakeCache.txt` as `ON` before either configure hits its (pre-existing,
environment-only) missing-dependency errors unrelated to this change.

## Testing Checklist

- [x] CMake configure (hipDNN) - `cmake -S projects/hipdnn -B build
-DHIPDNN_ENABLE_KERNEL_INGESTOR=ON` - Status: Passed (option registers
`ON` in cache)
- [x] CMake configure (hip-kernel-provider) - `cmake -S
dnn-providers/hip-kernel-provider -B build
-DHIPDNN_ENABLE_KERNEL_INGESTOR=ON` - Status: Passed (option registers
`ON` in cache)
- [ ] PR CI - GitHub PR checks - Status: Pending

## Technical Changes

- Adds `option(HIPDNN_ENABLE_KERNEL_INGESTOR "Enable the generic kernel
ingestor" OFF)` to `projects/hipdnn/CMakeLists.txt`, after the existing
option block.
- Adds the same option to
`dnn-providers/hip-kernel-provider/CMakeLists.txt`, alongside
`HIPKERNELPROVIDER_ENABLE_ROCKE`.
BrianHarrisonAMD added a commit to ROCm/rocm-libraries that referenced this pull request Aug 11, 2026
…-2402)

Build and test the ingestor on both superbuild jobs, mirroring how
HIPKERNELPROVIDER_ENABLE_ROCKE is enabled on the same configure lines.
The flag is declared by TheRock (ROCm/TheRock#7194) and defaults OFF, so
without this the engine is compiled out and gets no CI signal at all.

Both runners have a GPU, so this exercises the hipRTC compile and launch,
not just the build.
BrianHarrisonAMD added a commit to ROCm/rocm-libraries that referenced this pull request Aug 11, 2026
…-2402)

Build and test the ingestor on both superbuild jobs, mirroring how
HIPKERNELPROVIDER_ENABLE_ROCKE is enabled on the same configure lines.
The flag is declared by TheRock (ROCm/TheRock#7194) and defaults OFF, so
without this the engine is compiled out and gets no CI signal at all.

Both runners have a GPU, so this exercises the hipRTC compile and launch,
not just the build.
BrianHarrisonAMD added a commit to ROCm/rocm-libraries that referenced this pull request Aug 12, 2026
…-2402)

Build and test the ingestor on both superbuild jobs, mirroring how
HIPKERNELPROVIDER_ENABLE_ROCKE is enabled on the same configure lines.
The flag is declared by TheRock (ROCm/TheRock#7194) and defaults OFF, so
without this the engine is compiled out and gets no CI signal at all.

Both runners have a GPU, so this exercises the hipRTC compile and launch,
not just the build.
BrianHarrisonAMD added a commit to ROCm/rocm-libraries that referenced this pull request Aug 14, 2026
…-2402)

Build and test the ingestor on both superbuild jobs, mirroring how
HIPKERNELPROVIDER_ENABLE_ROCKE is enabled on the same configure lines.
The flag is declared by TheRock (ROCm/TheRock#7194) and defaults OFF, so
without this the engine is compiled out and gets no CI signal at all.

Both runners have a GPU, so this exercises the hipRTC compile and launch,
not just the build.
BrianHarrisonAMD added a commit to ROCm/rocm-libraries that referenced this pull request Aug 16, 2026
…ault-OFF flag (ALMIOPEN-2402) (#10606)

## Summary

Adds the foundation of the generic kernel ingestor ([RFC
0017](https://github.com/ROCm/rocm-libraries/blob/develop/projects/hipdnn/docs/rfcs/0017_UniversalKernelDescriptor.md)),
behind the default-OFF `HIPDNN_ENABLE_KERNEL_INGESTOR` flag. Every
descriptor kind (UED, KMD, UMD, UDD, UHD, UKD, KDP) exists as a real
type cross-referenced by id, and `KernelIngestorStateManager` answers
the three questions hipDNN's four host calls reduce to.

Two reference engines ship on it — `hipkernel:PointwiseAdd` and
`hipkernel:PointwiseSub` — described by descriptor data plus native
symbols resolved by name. Nothing is loaded from a file yet:
ALMIOPEN-2401 replaces one function body with a directory scan and the
rest keeps working.

JIRA ID : ALMIOPEN-2402

## What a second engine cost

`PointwiseSub` was added last, after the seam was finished, to measure
the claim. Three new files plus **12 lines across three shared files** —
one row in `IngestorPacks.cpp`'s table, two declarations in its header,
two CMake lines, one kernel-embedding entry. No matcher, state manager,
container, or engine factory changed.

A pack is two `.cpp` files: `<Op>Native.cpp` (graph matcher, kernel
matcher, scorer, dispatch handler, registration function) and
`<Op>Descriptors.cpp` (the descriptor set). No per-pack header.

The test side cost more — the shared graph-builder header became
pack-parametric and three test files followed — so "two files" describes
the production surface only.

## Risk Assessment

**Risk level 2.** Every new file and include site is behind `#ifdef
HIPDNN_ENABLE_KERNEL_INGESTOR`, default OFF. Verified with the flag off:
428 provider tests pass, no ingestor test in any binary, no `ingestor/`
path installed, touching an ingestor source rebuilds nothing.

Two changes reach shared code: the new header-only SDK surface under
`plugin_sdk/include/hipdnn_plugin_sdk/ingestor/`, which is what a second
provider would adopt; and an additive arch-name constructor overload on
`KernelCompileOptions`.

## ASIC Coverage

**Passing PR CI is not sufficient on its own.** Every
numeric-correctness assertion is `SKIP_IF_NO_DEVICES()`-gated and both
superbuild CI runners are device-less, so CI produces **zero signal on
correctness**. Local and cluster runs are the only proof.

Coverage is **gfx90a (MI210)** locally and **gfx942 (MI300X)**, both
green. No full multi-arch sweep is required: nothing here is
architecture-dependent beyond a hipRTC compile of a trivial elementwise
kernel through the provider's existing `IKernelCompiler` path.

## Design

**Layout.** The SDK half is header-only, so another provider can adopt a
piece without the whole ingestor. No new engine or plugin ABI interface.
Provider half at `engines/kernel_ingestor_engine/`.

**Two mechanisms with opposite lifetimes.** Native symbol registration
is permanent — the escape hatch survives the loader. Descriptor
provenance is temporary scaffolding standing in for files.

**A pack declares its architectures.** `KernelDescriptorPack::arch` is
checked ahead of every matcher.

**Symbols resolve at construction, not first use.** A descriptor naming
a symbol this build does not ship would otherwise survive load and
`isApplicable` and throw at plan build, past the point applicability is
a binding promise. Resolving in the constructor makes it a load-time
failure naming the descriptor.

**Interfaces worth a review opinion.**

- `IKernelDispatchHandler` splits into `workspaceBytes`, `prepare`, and
`launch`: the workspace query arrives before a kernel is chosen, and a
plan outlives the call that built it.
- `KernelSource` is a tagged union, not two strings: the kind is the
adapter dispatch point. Only `EMBEDDED_SOURCE` has an adapter; any other
kind is a load error.
- `ingestor::DeviceProperties` replaces `hipDeviceProp_t` on
`IDeviceResolver` and `MatchContext`, making `$device.*` a closed
vocabulary instead of ~90 implied ones and keeping HIP out of a header
another provider adopts.

## Testing Summary

Plugin SDK and provider unit tests, a concurrency suite (eight threads,
also under ThreadSanitizer), frontend E2E against the CPU reference
executor, and a flag-off pass covering build, tests, install surface,
and incremental rebuild.

Nine behaviours are mutation-verified — the fix reverted and the test
confirmed to catch it — including the cache-thrash fix and both packs'
refusal guards.

The refusal matrices are the load-bearing pack-layer tests: an
under-specified refusal accepts a graph the kernel cannot serve, which
is silently wrong rather than a missed optimization.

One branch is deliberately untested and marked so:
`HandleDeviceResolver` falling back when `hipGetDevice` itself fails
needs a broken HIP runtime to reach.

## Testing Checklist

- [x] Plugin SDK unit tests - `ninja hipdnn_plugin_sdk_tests &&
./bin/hipdnn_plugin_sdk_tests` - Status: Passed (237 tests)
- [x] Plugin SDK tests under ThreadSanitizer - `clang++
-fsanitize=thread` build of the concurrency suite - Status: Passed (zero
race reports)
- [x] Provider unit tests - `ninja hip_kernel_provider_tests &&
./bin/hip_kernel_provider_tests` - Status: Passed (509 tests, 1
pre-existing arch skip)
- [x] Provider integration tests - `ninja
hip_kernel_provider_integration_tests &&
./bin/hip_kernel_provider_integration_tests` - ASICs: gfx90a - Status:
Passed (16 tests, process exit 0)
- [x] Flag-off build and tests - same targets configured without
`HIPDNN_ENABLE_KERNEL_INGESTOR` - Status: Passed (428 provider tests, no
ingestor test present in any binary)
- [x] Flag-off install surface - `cmake --install` with the flag off -
Status: Passed (0 `ingestor/` paths installed)
- [x] pre-commit - `pre-commit run --files $(git diff --name-only)` -
Status: Passed
- [x] Cluster verification - ASICs: gfx942 - Status: Passed (237 SDK,
510 provider, 16 integration, exit 0; No skips on gfx942: the arch-gated
`TestContainer.GetApplicableEngineIdsSdpaGraph` that skips on the gfx90a
dev box runs here.
- [ ] PR CI - GitHub PR checks - Status: Pending

## Technical Changes

- **New SDK surface** under
`plugin_sdk/include/hipdnn_plugin_sdk/ingestor/`: the seven descriptor
types, `KernelIngestorStateManager`,
`GenericEngine`/`GenericPlanBuilder`/`GenericPlan`, `NativeRegistry`,
`SymbolScope`, `IKernelHeuristic`, `IKernelDispatchHandler`,
`IDeviceResolver`, `DeviceProperties`, `MatchContext`, `Catalog`,
`LruCache`, `MakeEngine`.
- **Provider engine** at `engines/kernel_ingestor_engine/`: pack table,
symbol-registration sweep, `discoverDescriptorSets()`, run-time
engine-name registration, shared device resolver.
- **Two reference packs**, two `.cpp` files each, plus their kernels.
- **`Container`** loops over discovered descriptor sets instead of
naming an engine.
- **CI**: `hipdnn-superbuild-ci.yml` configures both jobs with the flag
ON. The flag is declared in TheRock (ROCm/TheRock#7194).
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Status: Done

Development

Successfully merging this pull request may close these issues.

2 participants