Repository navigation
feat(hipdnn): add the generic kernel ingestor foundation behind a default-OFF flag (ALMIOPEN-2402) - #10606
Merged
BrianHarrisonAMD merged 42 commits intoAug 16, 2026
Conversation
BrianHarrisonAMD
changed the title
tvy-amd
reviewed
Aug 11, 2026
SamuelReeder
added a commit
that referenced
this pull request
Aug 11, 2026
BrianHarrisonAMD
force-pushed
the
users/bharriso/almiopen-2402-ingestor-skeleton
branch
from
August 11, 2026 21:09
cdaf79d to
a31904e
Compare
4 tasks
8 of 9 tasks
15 of 16 tasks
BrianHarrisonAMD
force-pushed
the
users/bharriso/almiopen-2402-ingestor-skeleton
branch
from
August 12, 2026 22:28
4576b09 to
4b18816
Compare
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
SamuelReeder
force-pushed
the
users/bharriso/almiopen-2402-ingestor-skeleton
branch
from
August 13, 2026 14:27
da6f632 to
6c30b18
Compare
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #10606 +/- ##
========================================
Coverage 69.29% 69.29%
========================================
Files 2777 2777
Lines 457793 457793
Branches 67512 67512
========================================
Hits 317194 317194
Misses 117467 117467
Partials 23132 23132
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
BrianHarrisonAMD
marked this pull request as ready for review
August 13, 2026 23:29
A pack whose kernel-scoped matchers rejected every kernel logged "admitted 0 of N kernel(s) after kernel-scoped matching", reading as a participant that happened to score zero rather than a pack that is out. The two declines above it, arch and graph-scoped matcher, both say plainly that the pack is excluded. Give zero admissions its own message and an explicit continue. The continue is a no-op today, being the last statement in the loop body; it is there so the skip sits where the decision is made and so a future per-pack step inherits it. mergeBound stays ahead of the admission check deliberately. It is the only detector for two matchers binding one token to disagreeing values, which is a descriptor authoring bug, invalid whether or not the pack admits a kernel on this particular graph. Deferring it would make detection depend on the workload under match. No behaviour a caller can observe changes.
The ingestor test fixtures leaked native-registry state between tests, so the suite only passed in declaration order. Under --gtest_shuffle one seed produced 147 failures, and most individual tests failed under --gtest_filter because they depended on a symbol another test registered. Three fixture bugs, all introduced with these fixtures in this PR: - makeStateManager() named "test.dispatch" without registering it, unlike makeTestDispatches(). It worked only once some other test in the binary had registered the symbol. - ensureNoopDispatchRegistered() documented itself as idempotent but used replaceSymbol(), clobbering a handler a test had installed through ScopedDispatchRegistration. It now registers only when the symbol is absent, which is what the comment always claimed. - makeStubDispatches() registered through a one-shot function-local static. A ScopedDispatchRegistration constructed before its first call captures a null previous entry and unregisters on destruction, so the static never re-ran and every later caller saw an unresolved symbol. That cascade was the source of the remaining shuffle failures. Two tests also unregistered by hand after their assertions, which leaks whenever the body throws, feeding the same cascade. They now hold the shared ScopedResource<T> guard instead, via a scopedGraphMatcher() helper beside the existing ScopedBlockSizeScore. Verified: 30/30 shuffle seeds green, and all 239 tests pass individually under --gtest_filter.
The skip added in af9a0c5 had no test. The log is the whole point of that change, and SharedLogRecorder makes it assertable, so cover it rather than leave the branch to a one-off manual check. Two tests: a pack whose kernel matchers reject everything must say it admitted no kernel, and must not fall through to the "admitted 0 of N" contributor message; a pack that admits some still reports the count. Mutation-verified: with the skip reverted the zero-admit test fails and the contributor test still passes.
BrianHarrisonAMD
force-pushed
the
users/bharriso/almiopen-2402-ingestor-skeleton
branch
from
August 14, 2026 16:31
b3f0a09 to
1d02966
Compare
jdcampbe
reviewed
Aug 14, 2026
jdcampbe
approved these changes
Aug 14, 2026
Merged
1 of 3 tasks
…(ALMIOPEN-2402) registerEngineName() throws on a hash collision, and it was called from the initializer of getEngineDefinitions()'s function-local static. An escaping throw costs HIP_MLOPS and ASM_SDPA their rows even though neither is implicated, and C++ leaves the static uninitialized, so the next call reruns the whole initializer and throws again -- measured with a probe, not assumed: the initializer ran on all three of three calls after throwing on the first. The try/catch is per set rather than around the loop, so a collision costs only the set that caused it and every other engine still registers. Reported in review. Samuel's ALMIOPEN-2401 removes provider-side registration entirely (the loader interns the name and this becomes engineNameToId), but the fix stands on its own until that lands.
…state (ALMIOPEN-2402) EngineRegistrar's constructor writes the name into the process-wide registry and the object has no members, so s_registrars was keeping alive something with nothing to keep. Nothing read it back -- it was the only reference to the type in the repo. It arrived alongside the string-interning fix and got carried along with it; the interning is what was load-bearing. s_names stays: the registry holds string_view, so a name built at run time still needs storage outliving it, and it stays a deque because reallocation would move a registered view's string. Checked bugprone-unused-raii, which is enabled in both .clang-tidy files: it does not fire on this, since EngineRegistrar has no destructor to discard. Reported in review.
…g again (ALMIOPEN-2402) GenericEngine::initializeExecutionContext populates the settings from engineConfig and stores them on the context immediately before calling buildPlan, which then parsed the same config a second time. Beyond the wasted work it was a second place for the two paths to disagree: getMaxWorkspaceSize already reads the stored filter, so the two knob consumers derived the same answer by two routes. buildPlan now reads executionContext.executionSettings() and engineConfig is unused there, matching SdpaBwdPlanBuilder. That turns TContext into a real requirement, so it is documented on the class alongside the TSettings one. The two test stub contexts had no-op setExecutionSettings and no accessor, which is why nothing caught this. Both now store the settings, and the two tests that called buildPlan directly run the sequence the engine runs -- initializeExecutionSettings, setExecutionSettings, buildPlan -- instead of a state the engine never produces. Mutation-verified: dropping the setExecutionSettings call makes HonorsAnExplicitKnobSettingOverTheHeuristicDefault fail, so the knob path is genuinely covered rather than passing incidentally. Reported in review.
…der (ALMIOPEN-2402) A UHD was mandatory: DescriptorSet::heuristic was a value, and the state manager rejected a null scorer outright. An engine with no ranking model yet had nothing to supply, so the skeleton's two packs each invented one. heuristic and EngineDescriptor::heuristicId are now optional. When absent, makeKernelHeuristic returns DeclaredOrderKernelHeuristic and logs a warning naming the engine. The fallback adds no ordering rule of its own; it scores every kernel alike and lets rank()'s existing tie-break decide, which is already priority descending then descriptor id ascending. So "no heuristic" means declared order, and ranking stays total, stable across runs, and deterministic -- the absence of a model costs selection quality, never reproducibility. A fallback that ranked on kernel metadata would instead outrank priority, which is the one signal an engine without a model still has; the tests pin that by giving the losing kernel the larger block size, so a metadata-sensitive scorer fails them. The warning is the reason this is not silent. An engine that means to declare its order and one still waiting on a model are otherwise indistinguishable from the outside, which is what jdcampbe asked for in review. It names the engine, since an unnamed warning cannot tell an operator where to look; makeEngine therefore describes the engine before moving it out of the set, or the name would be read from a moved-from string. The state manager's null-heuristic rejection stays. Absent and null are different: absent is a descriptor that declined to name a model, null is a caller that failed to build one, and only the first is legitimate. Verified: 246 plugin SDK, 509 provider (1 pre-existing arch skip), 16 integration on gfx90a, 428 flag-off. Mutation-verified: the warning test fails when the message stops naming the engine, and all three ordering tests fail when the fallback scores on kernel metadata instead of a constant.
…ph (ALMIOPEN-2402) Three review nits in the state manager, all on the same path. buildCatalog called completeMetadata() for every kernel of every pack on every uncached graph, rebuilding a KernelDefinition whose every field -- id, pack, dispatch, source, completed metadata, priority -- is context-independent. The definitions are now built once in validateAndIndexPacks, which was already computing the same completed tuples to check uniqueness and then discarding them, and buildCatalog copies them. Held parallel to _packs rather than keyed by kernel id, since nothing validates that ids are unique within an engine and a map would have quietly assumed it. The duplicate-tuple check scanned a vector, so validation was quadratic in an engine's kernel count. It is a std::set now; MetadataValues is an ordered map, so it already orders and no comparator is needed. The catalog cache hit returned *cached, copying a Catalog that LruCache::get had already copied out. It moves out of that local instead. Mutation-verified: disabling the duplicate-tuple guard fails RejectsTwoKernelsSharingAMetadataTuple, so the set conversion is covered rather than merely compiling. The integration suite is the check that matters for the precompute, since it runs real dispatch through the matcher path: 16 passed on gfx90a, alongside 246 plugin SDK, 509 provider, 428 flag-off.
…iopen-2402-ingestor-skeleton # Conflicts: # .github/workflows/hipdnn-superbuild-ci.yml
cderb
approved these changes
Aug 14, 2026
…2402) score() is supplied by the pack and resolved by symbol, so its value is outside rank()'s control and NaN is reachable. NaN compares false against everything, so it read as equivalent to every kernel while real scores stayed ordered among themselves -- not a strict weak ordering, which stable_sort requires, and therefore undefined behaviour. The visible damage is to the other kernels, not the NaN one: a probe of the exact comparator over 3 NaN scores reports 60 transitivity violations and sorts 7 after 1, so the worst kernel can be selected first. libstdc++ tolerates it; MSVC's iterator debug level asserts on it. Map non-finite to -infinity as the score enters the sort. The order stays total, NaN-scoring kernels sink below every finite one, and the existing priority -> kernelId tie-break still orders them, so a pack returning NaN loses selection quality without costing determinism. Infinities are already well-ordered and pass through untouched. Mutation-verified: with the normalization reverted, all three NaN tests fail and the infinity test correctly still passes. Note the determinism-only assertion this started as was too weak to keep -- stable_sort against a broken comparator is reliably wrong rather than random, so ranking twice and comparing passes even with the bug present. The test now asserts the finite subsequence is still descending, which is the invariant that breaks.
…IOPEN-2402) Every ingestor header is wholly inside #ifdef HIPDNN_ENABLE_KERNEL_INGESTOR, and all 16 are installed whenever the flag is set. The exported hipdnn_plugin_sdk INTERFACE target carried include directories and link dependencies but no compile definitions, so a consumer doing find_package(hipdnn_plugin_sdk CONFIG REQUIRED) received the headers on disk and they preprocessed to nothing -- the types simply absent, with no diagnostic naming the cause. In-tree this was masked by two side doors: add_compile_definitions scoped to projects/hipdnn, and a target_compile_definitions on the provider's object library. Neither reaches an out-of-tree consumer, and samples/example_engine_plugin is exactly that consumer. INTERFACE propagates to both the build-tree export and the installed hipdnn_plugin_sdkTargets.cmake, so in-tree and out-of-tree agree. Verified against a real install tree: a TU naming ingestor::Catalog through find_package compiles and runs. Mutation-verified by stripping INTERFACE_COMPILE_DEFINITIONS from the installed targets file, which reproduces the reported failure exactly -- "use of undeclared identifier 'hipdnn_plugin_sdk'". Flag-off still installs 0 ingestor paths and exports the definition 0 times.
hipdnn_backend_logging_shutdown_tests fails intermittently on the Windows superbuild with exit code 0xC0000409. That code reads as STATUS_STACK_BUFFER_OVERRUN but is misleading: MSVC routes std::abort() through __fastfail, and the abort here is the test's own, taken when a callback worker ends with zero invocations. The test asserted a guarantee the library never made. It documents null-before-drain as the intended shutdown contract, and dropping queued callbacks during static destruction is correct -- the user's handle may already be gone. But the readiness gate claimed it ensured callbacks "have been invoked at least once" while only ever proving they were ENQUEUED: it signalled as soon as the API call returned. Delivery happens later, on the shared logger's single background thread, behind two callbacks that sleep 1ms per message. So main() could return -- starting shutdown, which disables every callback before draining the queue -- while a worker's only message was still queued. That message is then dropped and the worker ends at zero. Signal readiness only once the callback has actually fired. Plain workers register no callback and keep today's enqueue-only behaviour, which is correct for them. Reproduced deterministically by widening the slow callback 1ms -> 40ms and pinning to 2 CPUs: 15/15 failures before, 15/15 passes after, with the same victim distribution as CI (later-registered workers starve first). With the real 1ms sleep restored, 40/40 pass. No production code changes, so the library's shutdown semantics are untouched. Whether loggerShutdown() should drain the queue before nulling the holders rather than after is a separate question about logging semantics at exit, and is deliberately not answered here.
BrianHarrisonAMD
deleted the
users/bharriso/almiopen-2402-ingestor-skeleton
branch
August 16, 2026 22:51
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Adds the foundation of the generic kernel ingestor (RFC 0017), behind the default-OFF
HIPDNN_ENABLE_KERNEL_INGESTORflag. Every descriptor kind (UED, KMD, UMD, UDD, UHD, UKD, KDP) exists as a real type cross-referenced by id, andKernelIngestorStateManageranswers the three questions hipDNN's four host calls reduce to.Two reference engines ship on it —
hipkernel:PointwiseAddandhipkernel: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
PointwiseSubwas added last, after the seam was finished, to measure the claim. Three new files plus 12 lines across three shared files — one row inIngestorPacks.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
.cppfiles:<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, noingestor/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 onKernelCompileOptions.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
IKernelCompilerpath.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::archis 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
isApplicableand 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.
IKernelDispatchHandlersplits intoworkspaceBytes,prepare, andlaunch: the workspace query arrives before a kernel is chosen, and a plan outlives the call that built it.KernelSourceis a tagged union, not two strings: the kind is the adapter dispatch point. OnlyEMBEDDED_SOURCEhas an adapter; any other kind is a load error.ingestor::DevicePropertiesreplaceshipDeviceProp_tonIDeviceResolverandMatchContext, 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:
HandleDeviceResolverfalling back whenhipGetDeviceitself fails needs a broken HIP runtime to reach.Testing Checklist
ninja hipdnn_plugin_sdk_tests && ./bin/hipdnn_plugin_sdk_tests- Status: Passed (237 tests)clang++ -fsanitize=threadbuild of the concurrency suite - Status: Passed (zero race reports)ninja hip_kernel_provider_tests && ./bin/hip_kernel_provider_tests- Status: Passed (509 tests, 1 pre-existing arch skip)ninja hip_kernel_provider_integration_tests && ./bin/hip_kernel_provider_integration_tests- ASICs: gfx90a - Status: Passed (16 tests, process exit 0)HIPDNN_ENABLE_KERNEL_INGESTOR- Status: Passed (428 provider tests, no ingestor test present in any binary)cmake --installwith the flag off - Status: Passed (0ingestor/paths installed)pre-commit run --files $(git diff --name-only)- Status: PassedTestContainer.GetApplicableEngineIdsSdpaGraphthat skips on the gfx90a dev box runs here.Technical Changes
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.engines/kernel_ingestor_engine/: pack table, symbol-registration sweep,discoverDescriptorSets(), run-time engine-name registration, shared device resolver..cppfiles each, plus their kernels.Containerloops over discovered descriptor sets instead of naming an engine.hipdnn-superbuild-ci.ymlconfigures both jobs with the flag ON. The flag is declared in TheRock (feat(flags): add HIPDNN_ENABLE_KERNEL_INGESTOR build flag TheRock#7194).