Repository navigation
feat(hipdnn): RFC 0018 UMD graph matcher and op-schema registry codegen - #10341
Closed
DarylHawkinsAMD wants to merge 61 commits into
Closed
DarylHawkinsAMD wants to merge 61 commits into
DarylHawkinsAMD wants to merge 61 commits into
Conversation
DarylHawkinsAMD
force-pushed
the
users/dahawkin/rfc-umd-match-descriptor
branch
from
August 4, 2026 17:13
1989a8a to
f1578d2
Compare
jdcampbe
added a commit
that referenced
this pull request
Aug 5, 2026
- Add cderb as a contributor. - Cite the UMD RFC's working implementation (PR #10341): its JsonLogic.hpp compile-once/evaluate-many Expression is the concrete evaluator the feature extractor reuses; note `shape` is a compiler short-hand (not a runtime op) and the custom-operation hook is deferred in the PoC (the boundary where custom_library takes over). - Restore the two RFC-0017-aligned Section 10 corrections the restructure reverted: model cache is process-scoped (not per handle), and lazy load is triggered by ranking (not winning). - Add Section 10.3 "Efficient evaluation": lower features_signature at load, shared-prefix/per-candidate-suffix rows, reuse matcher bindings, single-candidate short-circuit -- the "expressive spec, fast hot path" guidance, grounded in the PoC's compiled Expression. Co-Authored-By: Claude <noreply@anthropic.com>
4 tasks
…ph-uuid-identity # Conflicts: # projects/hipdnn/backend/src/PlatformUtils.windows.cpp # projects/hipdnn/backend/tests/descriptors/TestGraphDescriptor.cpp # projects/hipdnn/flatbuffers_sdk/include/hipdnn_flatbuffers_sdk/utilities/json/Graph.hpp
Base automatically changed from
users-bharriso-rfc-ukd-kernel-ingestor
to
develop
August 5, 2026 23:32
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
jdcampbe
added a commit
that referenced
this pull request
Aug 6, 2026
…14.6) Add a generation-pipeline subsection on deriving a first-pass features_signature from what a package already carries, framed as two layers: Layer 1 op-intrinsic vocabulary (tensor/dim bindings, device props, and now arithmetic intensity) that is shared across all packages of an op and emitted by the UMD op-schema registry; Layer 2 the per-UHD features_signature that selects among them. Specify arithmetic intensity (FLOPs/byte) and what authors must provide: per-op umd_flops / umd_bytes table-level .fbs annotations, declared in data_types.fbs and applied like umd_opcode, reusing the PR #10341 codegen. Cover the mixed-dtype byte-formula caveat, the native-predicate fallback for non-closed-form ops, the superset/pruning caveat, and a data-free structural first pass (tile-fit ranking) as the zero-benchmark start. Extend Open Question 4 with the annotation dependency. Co-Authored-By: Claude <noreply@anthropic.com>
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## develop #10341 +/- ##
===========================================
+ Coverage 68.90% 69.16% +0.26%
===========================================
Files 2766 2765 -1
Lines 455513 455350 -163
Branches 67144 67134 -10
===========================================
+ Hits 313834 314914 +1080
+ Misses 118765 117555 -1210
+ Partials 22914 22881 -33
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
|
🎉 All checks passed! This PR is ready for review. |
Adds the structural skeleton of the RFC 0017 kernel ingestor, behind the default-OFF HIPDNN_ENABLE_KERNEL_INGESTOR flag, proven end to end for one hardcoded case: a single-node pointwise ADD over 1-element tensors. Every descriptor kind exists as a real, cross-referenced-by-id type (UED/KMD/UMD/UDD/UHD/UKD/KDP), and KernelIngestorStateManager answers the three questions hipDNN's four host calls reduce to. Nothing is loaded from a file: each UMD, UHD, and UDD names a native symbol resolved from a registry, which is the escape hatch RFC 0017 already defines for steps its declarative form cannot express. Replacing a native symbol with real descriptor data is a change to one descriptor field, not to the ingestor -- which is the property this arrangement exists to demonstrate. The generic machinery is header-only in plugin_sdk so other providers can adopt a piece without adopting the whole ingestor. It is excluded from the SDK install tree while the flag is off, so a default build's installed artifacts are unchanged. Dispatch performs a real hiprtc compile and launch rather than a recorded no-op: the parts most likely to be wrong when real kernels arrive are the compile path and uid-to-pointer argument resolution, and only a real launch exercises them. The provider's engine table gains one static row. Engine registration stays compile-time; runtime loading of engine descriptors (ALMIOPEN-2401) replaces that row with one built per installed UED file. Testing: - hipdnn_plugin_sdk_tests: 26 new tests covering the catalog cache, the registry's fail-closed lookup, ranking and its tie-break order, and the state manager's matching, pruning, caching, and validation. - hip_kernel_provider_tests: 31 new tests covering the descriptor set, both matcher shapes, workspace sizing, and a real GPU launch (gfx90a). - hip_kernel_provider_integration_tests: 9 new tests driving the plugin's C ABI through dlopen -- applicability, knob reporting, workspace, plan build, execute -- asserting c == a + b on device. - Flag off: 428 unit + 5 integration pass and the ingestor tests are absent, confirming the gate compiles the feature out entirely. JIRA ID : ALMIOPEN-2402
…MIOPEN-2402) Two defects in the kernel ingestor skeleton, both found by asking what happens when several handles share one provider container. The catalog cache is keyed on (graph, device), but the device was captured once when the engine was built rather than read from the handle each call. One container -- and therefore one cache -- is shared by every handle in the process, so two handles on different devices would have shared a catalog entry belonging to whichever device happened to be current at construction. That is a wrong answer, not a missed cache hit. IDeviceResolver now resolves the device per call from the handle's stream, and the dispatch handler compiles against the properties the call carries instead of ones captured at construction. The dispatch handler was owned by the engine but registered in a process-wide registry. A container is destroyed when its last handle closes, so the registration outlived the handler it pointed at, and the next container resolved freed memory. The handler and its compiler are now process-lifetime, matching the registry that holds them. The second defect is worth describing precisely, because it hides well. Every test passed with it present: a freed handler carrying no per-instance state keeps answering through its stale vtable. It surfaced only as a segfault at process teardown, after the last test had already reported success. Verified by mutation: reintroducing engine ownership makes the integration binary exit 139 on every run, and the fix returns it to 0. Testing: - 8 new concurrency tests: the LRU under concurrent readers, writers and eviction; the state manager serving one graph, several devices, uncacheable graphs and constant eviction from 8 threads; concurrent registry resolution. - Clean under ThreadSanitizer: 8/8 pass, zero race reports. - A regression test that the registered dispatch still resolves and launches after a container is destroyed and rebuilt. It documents what it cannot prove: the use-after-free is caught by the integration binary's exit status, not by any assertion. - Full suites green: 125 plugin_sdk, 460 provider unit, 14 integration, all exiting 0. Flag off unchanged at 428 unit + 5 integration. JIRA ID : ALMIOPEN-2402
The input/output tensor bindings looked up a tensor by a field called 'role'; rename it to 'name' for consistency with AttrBinding.name and the umd_name annotation it carries. Updates the neutral types, the generator (bindName local, seenNames, diagnostics) and regenerated header, and the tests. (cherry picked from commit d0aca81) (cherry picked from commit 46c5ea5)
Align the docs with the code rename: a tensor's lookup key is a 'name', not a 'role'. Rewords the concept throughout (op-schema names, name-map/name-bind grammar, A.3 'Names and edges', B.5 registry example keys) and lowercases the q/k/v/o example strings to match the schema. (cherry picked from commit 27fcafd) (cherry picked from commit af207fe)
OpSchemaEntry stored the union discriminant as int, forcing a hard static_cast<int> at every generated entry and an inverse cast in lookupOpByType. Carry the real NodeAttributes enum value instead: add an opaque forward declaration (fixed uint8_t underlying type) of the enum to the neutral OpSchemaRegistry.hpp so the field can hold the enum without pulling the generated graph schema into that header. Both casts are gone; lookupOpByType takes and compares the enum directly. The generated registry TU sees both the forward declaration and the real definition, so any underlying-type mismatch is a compile error there. (cherry picked from commit 485d227) (cherry picked from commit 6f4f5da)
… findings
The umd op-schema registry header is machine-generated (umd_registry_gen)
and lives under umd/, outside the data_objects path the clang-tidy
ExcludeHeaderFilterRegex already skipped, so its emitter-artifact
diagnostics (C-style arrays, non-conforming identifiers, brace-less
statements) were being reported. Extend the exclusion to umd/*_generated.
{h,hpp}, scoped to that directory rather than a blanket *_generated so a
hand-written file elsewhere is still checked.
Fix the remaining hand-written findings the audit surfaced:
- umd_registry_gen.cpp: brace all statements; make pointer/size bool
conversions explicit; return a braced init list; rename kNs -> NS and
static local k -> s_keywords; replace chained-`+` fail() messages with a
failField() helper that appends.
- OpSchemaRegistry.hpp: rename AttrType constants Int/Float/Bool/Dtype to
UPPER_CASE (INT/FLOAT/BOOL/DTYPE) and update the defaults and comments.
- TestOpSchemaRegistry.cpp: brace single-statement if/for, make range-loop
name const, and use the UPPER_CASE AttrType constants.
- op_schema_registry_generated.hpp: regenerated from the updated generator.
clang-tidy reports zero errors on both changed TUs; the flatbuffers_sdk
tests build and UmdOpSchemaRegistry.* pass.
(cherry picked from commit 616591f)
(cherry picked from commit 550915b)
Add the provider-local Universal Match Descriptor runtime under dnn-providers/hip-kernel-provider/include/hip_kernel_provider_common/umd: - BindingContext: the queryable bindings object and JsonLogic data source, resolving the RFC's tensor/graph/attribute/device namespaces over a live flatbuffer graph via the Phase 0 generated registry readers plus a hand-written tensor-path resolver. Fails closed on any unresolved reference. - UmdCompiler: parses and validates a hipdnn.umd/v1 descriptor (Appendix A.10), lowers the shape short-hand and layout aliases, statically type-checks the criteria against the generated AttrType, and compiles it to a jlogic::Expression. Rejects malformed input with UmdCompileError. - UniversalGraphMatcher: owns exactly one compiled descriptor (1:1) and matches a graph via structural subgraph binding with edge consistency across shared pattern variables, then evaluates the criteria. Bounded backtracking. - UmdPathParse: shared variable-path parsing used by both resolvers. Also lowercase the operand/result name-map keys in the RFC 0018 examples to match Appendix A.3 and the generated registry. Exercised by TestUniversalGraphMatcher (35 cases): compiler validation, single/multi-node structural matching, edge consistency, per-node attributes, namespace resolution, and fail-closed behavior. (cherry picked from commit f00bcf9) (cherry picked from commit 7053605)
Resolve `$kernel.<field>` criteria references against caller-supplied, fully-resolved KMD metadata. Kernel fields are DYNAMIC scalars: not statically typed at UMD-compile time (their types come from the engine's KMD, validated at load), unified with any scalar domain at match time. - JsonLogic: add referencesVariableRoot() first-path-token helper. - UmdCompiler: DYNAMIC kind for $kernel fields; flag descriptors that reference kernel metadata; accept DYNAMIC in bool/numeric/if positions; restrict DYNAMIC unification to scalar domains (reject array/tensor). - BindingContext: bindKernelMetadata() backed by a JsonDataSource. - UniversalGraphMatcher: three-arg match() overload taking kernel metadata; two-arg overload throws for kernel-referencing descriptors; referencesKernelMetadata() routes between them. - Tests for helper, compiler type domain, and matcher routing/behavior. (cherry picked from commit 1607859) (cherry picked from commit 2cd003d)
The CMake build regenerates op_schema_registry_generated.hpp in-tree from graph.bfbs, but the flatc-hipdnn pre-commit hook only refreshes the --cpp data-object headers, so a committed schema change can leave the registry stale. Add a superbuild CI step that fails if the post-build tree is dirty at that header, and document the split of responsibility in the hook config and run_flatc.py. (cherry picked from commit 8b16f55) (cherry picked from commit 1d6c72f)
RFC 0017 grew from 1318 to 2701 lines and renumbered every section from §8 onward, breaking 13 cross-references and tightening several rules the UMD spec restates. Cross-references: retarget all 13 broken anchors (old §8.3 -> §9.3, §9 -> §10, §10 -> §12, §12.x -> §14.x, §14 -> §16, §15 -> §17) and fix a pre-existing label/anchor mismatch that pointed "§9" at the packaging section. Contract changes: - A UMD carries `version` and `sdk_version`, gated against a runtime floor and compared numerically by (major, minor). - Add the `present` / `not_present` operators, which always evaluate rather than declining on an absent reference, and use the list form to decline unsupported optional operands in one call. - Widen `value_or_default` so the fallback may be any type-compatible expression, including a second field reference. - Add `$q.is_runtime_pass_by_value`, `$q.value_f32`, and `$graph.is_override_shape_enabled`, and describe precomputed fields as a distinct mechanism between the built-ins and the escape hatch. - State the exactness-to-kernel rule (every baked quantity is a KMD field pinned by a `$kernel.*` criterion) and the umbrella-matcher rule, both marked as pack-level loader checks rather than per-descriptor ones. - Defer matching order and caching to RFC 0017 §8: matching runs during applicability against a provider-owned cache keyed on engine, graph, and device, producing a catalog and bound token state. Cite the base-path invariant as a constraint on match semantics. - Narrow day-one fusion matching to a pattern that is the entire graph. - Memoize `$kernel.*` matchers on the fields they read, not on the whole metadata tuple. Layout aliases now document the `extractStrideOrder` encoding the codebase computes (`nchw` -> [3,2,1,0]) instead of the inverse position-to-dim convention, which no implementation used. Retarget observability at §10 and the tooling claim at the new §11, and record the runtime opt-out env vars and their shared-matcher risk. Refresh all three diagrams: JsonLogic syntax instead of the pre-JsonLogic key/value form, per-tensor dim names, `$sdpa_fwd.*` instead of `$root.*`, lowercase operand keys, and the applicability-time cache. (cherry picked from commit 7f061b2) (cherry picked from commit 4a47215)
RFC 0017 tightened the match contract on the base branch. Five changes were load-bearing; the sixth replaces a bool the new rules cannot use. The blocking one is the version keys. `validateTopLevel` rejects any unknown top-level key, so a spec-conformant descriptor carrying `sdk_version` failed to compile at all. Both `version` and `sdk_version` are now accepted, optional so existing descriptors keep compiling, and gated against a floor. Comparison is numeric on (major, minor), so "1.10" ranks above "1.9"; anything that does not parse is below every floor. `present` / `not_present` are new operators. Unlike every other operator they do not propagate the null an unresolved path reads: answering "was this supplied?" always yields a boolean, which is what lets a pack write "absent, or present and constrained" instead of silently accepting a graph it cannot serve. Both fold with `and` over their arguments, so one call declines a whole list of unsupported optional operands. They resolve the bare reference rather than the `.present` field path, so they work on required operands too; the field keeps its existing optional-only rule. `value_or_default` now type-checks both arms and accepts an expression or a second field reference as the fallback, not just a literal. Three namespace fields join the vocabulary: `$graph.is_override_shape_enabled` (the graph's own opt-in, distinct from the descriptor's `allow_override_shape`), `$q.is_runtime_pass_by_value`, and `$q.value_f32`, which coerces whichever arm of the tensor's eight-arm value union is set. A tensor carrying no compile-time value resolves to null so a criterion over it declines, rather than reading a zero that would satisfy a comparison. The Float8Value arm decodes its raw bits through the tensor's data_type; a pairing that names no 8-bit format also resolves to null. `referencesKernelMetadata` becomes the set of `$kernel.*` field names the criteria read. A matcher is memoized per distinct value of the fields it reads, not per kernel, and a loader must check each name exists in the engine's KMD -- neither is expressible with a bool. The set and the bound-symbol set are both exposed on the matcher so a KDP/KMD loader can consume them. Two existing tests now assert on the field set instead of the bool. (cherry picked from commit f1578d2) (cherry picked from commit a24c426)
…sion gating
Three RFC 0017 contracts the matcher did not honor. The first is a
fail-open bug in the applicability path.
`present`/`not_present` were documented as special because they alone do
not propagate the null an unresolved path reads. They were in fact the
only operators that handled null at all; every other one coerced it, so
a field read on an ABSENT optional operand read as 0 / false /
not-equal:
{"!": "$bias.packed"} => true
{"!=": ["$bias.dtype", "BFLOAT16"]} => true
{"<": ["$bias.rank", 5]} => true
{"==": ["$bias.dtype", "$mask.dtype"]} => true (two absent "equal")
Every one of those is a narrowing check that PASSED on an operand the
pack cannot serve, so a matcher accepted graphs it had no kernel for.
RFC 0017 §5 requires such a check to "neither pass nor fail, it simply
does not run".
Null now means unresolved and propagates through every operator except
the two presence predicates and `value_or_default`, which answer "did
this resolve?" and so always yield a real value. `and`/`or` are Kleene:
a definite false still decides an `and` and a definite true an `or`,
which is what makes RFC 0017's own "absent, or present and constrained"
idiom work when the operand is missing. A null root is falsy, so an
undecided criterion declines. `/`, `%`, and `ceil_div` decline on a zero
divisor instead of yielding inf/NaN, and `log2`/`rsqrt` on a
non-positive argument.
Version gating was inverted. RFC 0017 §4 makes `version` a CEILING --
reject a differing major or a minor newer than the runtime's, and "a
file stamped 1.0 loads on a 1.1 runtime" -- but both keys were gated
against a floor, so a conformant older descriptor was refused by a newer
runtime, exactly backwards. `sdk_version` is a floor the GRAPH sets, not
a constant, and that half was missing entirely; hipDNN already computes
it as `min_required_engine_api_version` from the optional features a
graph uses, so the matcher reads that rather than deriving a second
source of truth. An omitted key now reads "1.0" instead of empty, so the
per-graph comparison always has a concrete version.
`$q.value_f32` dropped FP8_E8M0 -- an 8-bit format its own comment
claimed to cover -- to the fail-closed null, and decoded INT8 through
UINT8's unsigned widening, so a stored -1 read as 255.0.
Tests: the two that encoded the inverted floor now assert the ceiling
and the older-minor case. New coverage for the fail-open forms, the
Kleene short-circuits, the per-graph floor, and the seven value-union
arms plus the fp8 decoder, none of which had any.
(cherry picked from commit 0c6df1c)
(cherry picked from commit 7731a28)
…claims The prior alignment pass restated several RFC 0017 contracts in the direction the implementation had them, not the direction RFC 0017 states them. RFC 0017 is the authority and is left untouched here. Version gating. §13, A.1, and A.10 §1 all required `version` to be "at or above the runtime's floor". RFC 0017 §4 makes it a CEILING: reject a differing major or a minor newer than the runtime's, otherwise load -- "a file stamped 1.0 loads on a 1.1 runtime". Under the floor wording that descriptor was refused, exactly inverting the rule. `sdk_version` is a floor, but one the GRAPH sets per match, not a runtime constant; both bounds now appear, and the graph-side floor is identified as hipDNN's existing `min_required_engine_api_version` rather than a new mechanism. Absent-optional semantics. A.4 said a field read on an absent optional "declines the whole criteria". That is the outcome, not the mechanism, and it left `and`/`or` unspecified -- under a plain decline, RFC 0017 §5's "absent, or present and constrained" idiom cannot work, since the second arm's field reads never run. Stated as three-valued: the read resolves to unknown, unknown propagates rather than coercing, and a definite result still short-circuits either connective. A.7 carries the same rule for the operator table. Layout. The fused Conv-Bias-ReLU example named `$y`'s dims in NHWC order while its `stride_order` of [3,0,2,1] implies an NCHW logical order -- `shape` names logical dims, layout is the stride order. Corrected, and the claim that the alias set "matches the layouts validateSupportedLayout accepts today" narrowed: that oracle covers the four conv layouts, while `bhsd` and `contiguous` are additions for attention. `extractStrideOrder` is cited at its implementation rather than a section comment. RFC 0017 §5's `stride_order` literals use the inverse, position-to-dimension convention, which no implementation uses and which RFC 0017 never states as intended. Raised as Open Question 6 for RFC 0017 to resolve, alongside the existing Open Question 5, rather than changed here: a matcher copied from those examples silently accepts the wrong physical layout. The PoC plan expanded "nhwc" to [0,2,3,1], contradicting the A.8 table it cites and the encoding `extractStrideOrder` computes, so implementing to it would have baked in the wrong permutation. (cherry picked from commit 8e5b361) (cherry picked from commit dd12b10)
RFC 0018 and the matcher published `stride_order` as `extractStrideOrder` returns it: a per-dimension stride RANK, entry `d` giving dimension `d`'s rank with higher meaning slower-varying, so NHWC is [3,0,2,1]. RFC 0017 §5 writes the inverse: an ordered list of logical dimension indices, outermost first, so NHWC is [0,2,3,1]. The two forms are exact inverses and carry identical information -- the conversion round-trips over every permutation, and reproduces every RFC 0017 literal including §13.2's BSHD [0,2,1,3]. Neither document was wrong about layouts; they spelled the same layouts two ways. RFC 0017 is the authority, so RFC 0018 moves. Its form is also the better one to author in: [0,2,3,1] over an (n,c,h,w) dim order reads N,H,W,C left to right. The rank vector does not read as anything, and misreading one as the other is silent -- take NHWC's [3,0,2,1] as an outermost-first list and it names H,C,W,N, a layout that does not exist. Converting the spec alone would have been worse than leaving it: an author would write [0,2,3,1] while the matcher compared against [3,0,2,1], and the criterion would decline forever with nothing to show why. So BindingContext now inverts once, at the single point a descriptor-visible value is minted, and the A.8 alias table is expanded in the same form. Nothing in the graph model changes; extractStrideOrder and every existing caller are untouched. Cross-tensor equality (`$x.stride_order == $y.stride_order`) is convention-independent and reads the same either way. The new test pins the direction with an NHWC graph rather than a contiguous one: [0,1,2,3] and [3,2,1,0] are both palindromic under this inversion, so a contiguous tensor cannot tell a correct conversion from a missing one. Drops Open Question 6, which asked RFC 0017 to invert its literals. (cherry picked from commit 9ad934d) (cherry picked from commit cecfc9e)
Two independent CI build failures on the UMD matcher change. Windows: <minwindef.h> defines IN as an empty annotation macro, so the jsonlogic Op::IN enumerator expanded to nothing and JsonLogic.hpp failed to parse in any translation unit that also pulls in windows.h. TestUniversalGraphMatcher.cpp does, via the flatbuffers/test SDK headers, which is why it broke while TestJsonLogic.cpp did not. Rename the enumerator to Op::MEMBERSHIP; the JsonLogic operator key stays "in", so rules and the wire format are unchanged. Linux: the superbuild runs clang-tidy as part of the compile, and modernize-return-braced-init-list (warnings-as-errors) rejected graphRequiredSdkVersion's explicit std::string construction. Return a braced initializer list instead. (cherry picked from commit 1be18f1) (cherry picked from commit 03d1ce3)
The test-name validators require a suite to read (Test|Integration)[Gpu]FeatureName[Datatype]. The six suites added for JsonLogic and the UMD matcher omitted the Test prefix, so hipdnn_test_name_validation and hip-kernel-provider_test_name_validation reported 129 non-conforming names. These only surfaced now because the build failure upstream of them skipped the test step. Prefix the suites: JsonLogic, JsonDataSource, BindingContext, UmdCompiler, UniversalGraphMatcher, and UmdOpSchemaRegistry. Test case names are unchanged; the JsonLogic.md gtest_filter example follows the rename. (cherry picked from commit a993471) (cherry picked from commit 47cae9a)
develop added ResampleBwdAttributes and MoeGroupedMatmulAttributes to the NodeAttributes union. The registry enumerates that union, so the committed header was two entries short. Neither table carries umd_* annotations, so both land as attribute-only entries keyed on their table names. (cherry picked from commit 6e09bbe) (cherry picked from commit 5afc3f6)
…lude root AsmSdpaConfigHelpers.cpp reached for SdpaPlanUtils.hpp through "../../engines/...", which resolves relative to the file's own directory (src/tests/engines/asm_sdpa_engine/plans, nonexistent) and to no -I root on the command line, so the translation unit did not compile. ConfigHelpers.cpp beside it already spells the same header relative to the src include root; match it. That root is not universal, though: this file is compiled into two targets. hip_kernel_provider_tests links hip_kernel_provider_private, which exports src as a PUBLIC include directory, so the root arrives there for free. hip_kernel_provider_integration_tests pulls the same source in but links only the SDKs, so it inherits nothing from the provider and has to name the root itself. Without it, the respelling merely moves the failure from one target to the other -- and only in a configuration that builds the integration tests, which is why it survived a unit-test-only build. Added beside the codegen and shared-helper entries the same block already lists. (cherry picked from commit a6e693d) (cherry picked from commit 974237b)
RFC 0018 §8 put the native-predicate hatch inside `criteria`: a custom operation named by a namespaced key, resolved against a predicate registry the compiler consults, nesting and negating like any built-in. That buys nothing the ingestor does not already provide, and it costs the property that makes the criteria language auditable -- that A.7 is exhaustive. A compiler that resolves names against a runtime registry cannot validate a descriptor from the descriptor alone: the same file compiles or fails depending on which predicates the running provider happens to ship, and the registry becomes a second, undocumented half of the operator reference. The hatch now sits beside the descriptor rather than inside it. A check that needs real C++ is a native matcher: an ordinary `GraphMatcherFn` registered in the provider's `NativeRegistry` and named by a `MatchDescriptor`'s `matchSymbol` (RFC 0017 §5). A pack lists it alongside its descriptor-backed matcher and the ingestor conjoins the verdicts -- a pack's graph-scoped matchers all run and all must pass -- so "this UMD's criteria AND this C++ predicate" is two matcher ids, not one nested inside the other. The two SDPA gates that motivated the hatch (strides_fit_u32, sdpa_mask_consistent) move out of the example descriptor's criteria on that basis. The compiler follows: with no registry to resolve against, a namespaced key is not a special case. Dropping the dedicated rejection lets it fall through to the unrecognized-operator error every other unknown key takes, which is the same fail-closed outcome by the general rule instead of a bespoke one. `RejectsCustomOperation` becomes `RejectsUnknownOperator` and now asserts that, rather than asserting a PoC-scope message that described a feature no longer planned.
DarylHawkinsAMD
force-pushed
the
users/dahawkin/rfc-umd-match-descriptor
branch
from
August 12, 2026 17:25
934a2b3 to
78d07b4
Compare
A compiled descriptor kept its id as the descriptor's text UUID, while the plugin SDK's DescriptorId (Descriptors.hpp) is already a 128-bit hipdnn_flatbuffers_sdk::utilities::UuidBytes. Two representations of one identity means every future cross-reference between a matcher and the descriptor set that names it pays a parse. CompiledUmd::id, MatchResult::umdId, and UniversalGraphMatcher::umdId() now hold UuidBytes. The JSON wire form is unchanged: a descriptor still declares `"id": "<uuid>"`, and the compiler parses it once. The hand-rolled isUuid() string check is gone; parseUuid() from the SDK is now the single source of truth for the A.10 section 1 well-formedness rule, with its std::invalid_argument translated to UmdCompileError so the fail-closed contract is unchanged. The accepted id set is identical to before. The $kernel-metadata routing diagnostic formats via formatUuid(), so its text is unchanged.
BrianHarrisonAMD
force-pushed
the
users/bharriso/almiopen-2402-ingestor-skeleton
branch
from
August 12, 2026 22:28
4576b09 to
4b18816
Compare
SamuelReeder
force-pushed
the
users/bharriso/almiopen-2402-ingestor-skeleton
branch
from
August 13, 2026 14:27
da6f632 to
6c30b18
Compare
BrianHarrisonAMD
force-pushed
the
users/bharriso/almiopen-2402-ingestor-skeleton
branch
from
August 14, 2026 16:31
b3f0a09 to
1d02966
Compare
Base automatically changed from
users/bharriso/almiopen-2402-ingestor-skeleton
to
develop
August 16, 2026 22:51
Contributor
|
This pull request has been inactive for 25 days and will be marked as stale. If you would like to keep this PR open, please:
This PR will be automatically closed in 5 days if no further activity occurs. |
Contributor
|
This pull request has been automatically closed due to inactivity (30 days with no updates). If you'd like to continue working on this, feel free to reopen the PR or create a new one. |
jdcampbe
added a commit
that referenced
this pull request
Sep 28, 2026
JIRA ID : AICK-1699 Close date: August 21st ## Summary Adds **RFC 0019: Universal Heuristic Descriptor (UHD)** — the "UHD + kernel selection" follow-up named in RFC 0017 §12.2 ([PR #9533](#9533), open). It designs the data-driven, per-engine kernel-selection model for the descriptor-based provider, plus the tooling that generates one. Docs-only: one new file, `projects/hipdnn/docs/rfcs/0019_UniversalHeuristicDescriptor.md`. ## What it covers - **UHD schema & adapters** — one discriminant (`adapter`) over `static_order` / `table` / `tree_data` (default) / `onnx` / `custom_library`; the model ships **as data with the engine**, read by an in-tree interpreter (no `liblightgbm` runtime dependency). - **Ownership** — the UED (engine) owns one UHD and one KMD; a KDP joins the engine. Knobs are a view onto KMD fields; the KMD↔UHD coupling is conditional (only a *breaking* KMD change forces a retrain — additive changes and dispatch-only fields do not). - **Feature contract** — an inline `features_signature` in JsonLogic (the language pinned by the sibling UMD RFC), sourced from the matcher's bound symbols; a three-part fail-closed contract check. - **Efficient evaluation** — the `features_signature` lowers to a compiled expression at load (reusing the UMD PoC's compile-once `JsonLogic.hpp`), with a shared per-graph feature prefix and a per-candidate suffix, so the hot path stays tight while the data contract stays extensible. - **Engine-selection integration** — applicability bubbles up before ranking; the tooling produces **two** heuristics (a cheap engine-level performance estimate and the fine-grained config UHD) consumed by **two** policies (quick vs. thorough), with a rank-ordering fallback. Cardinal (calibrated-TFLOPS) scoring enables cross-engine comparison. - **Generation pipeline** — a standalone tool over hipDNN's public **autotune** API ([RFC 0013](https://github.com/ROCm/rocm-libraries/blob/develop/projects/hipdnn/docs/rfcs/0013_Autotune.md)); ships a trivial `static_order` pack, then generates and drops in a `tree_data` UHD — usable by any package author, no provider code changes. ## Relationship to concurrent RFCs - **Parent: RFC 0017 (Universal Kernel Descriptors)** — [PR #9533](#9533) (open into `develop`; not yet merged there). This RFC is aligned with the review-feedback updates that landed on the RFC 0017 branch via [PR #10039](#10039): knob-as-exposed-KMD-field, cache key `(engine id, graph id, device id)` + inventory generation counter, lazy-load-on-ranking (not on winning), and the conditional KMD↔UHD coupling / KMD-superset-of-features relationship. - **Sibling: RFC 0018 Universal Match Descriptor (UMD) & graph matcher** — [PR #10341](#10341), with a working implementation. Its `JsonLogic.hpp` (compile-once / evaluate-many, `$`-sigil convention) is the concrete evaluator this RFC's feature extractor reuses; this RFC also answers that RFC's open question by making the matcher's bindings the canonical feature source. **Numbering note:** the UMD RFC holds "0018", so this doc is **0019**; final numbers reconcile at review across the RFC 0017 §12.2 follow-up series. ## Test plan - [ ] Docs-only change — no build/test impact. - [ ] Markdown renders; internal section anchors resolve (audited: 99 refs, 0 broken). - [ ] RFC review by hipDNN descriptor-series owners. - [ ] Confirm the 0018/0019 numbering with the UMD RFC author (dahawkin, PR #10341). --------- Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Christopher Erb <christopher.erb@amd.com>
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
Implements Phase 1 of RFC 0018 (Universal Match Descriptor): a declarative graph matcher that decides whether a kernel applies to an incoming problem graph and, in the same pass, binds the named variables a dispatch/workspace formula references. This replaces the graph half of a hand-coded
IPlanBuilder::isApplicablewith data, so a family of near-identical kernels shares a handful of matchers instead of carrying a bespoke C++ check each.The symbol table is not hand-maintained: a build-time tool reads the FlatBuffers reflection schema and generates a strongly-typed op-schema registry from
umd_*annotations on the existing.fbsfiles, so the bindings stay in lockstep with the schema by construction. SDPA forward is annotated as the worked example.JIRA ID : AICK-1698
This PR is stacked. It targets #10606 (the generic kernel ingestor foundation, ALMIOPEN-2402), which in turn targets #10377 (persistent graph UUIDs). The expression layer here is the matching implementation that ALMIOPEN-2403 adopts to displace
MatchDescriptor::matchSymbol—Descriptors.hppin #10606 names this follow-up explicitly. It should merge after both.Risk Assessment
Medium risk (3). Nothing here is reachable from a shipping code path: the matcher is header-only and consumed exclusively by unit tests, and it is not registered as a provider engine. The real blast radius is build infrastructure —
HIPDNN_GENERATE_SDK_HEADERS(default ON) now also builds a host tool (umd_registry_gen) and regenerates a committed header in-tree on every configure/build, so a flatc or reflection-API hiccup surfaces as a build failure rather than a runtime one. Schema edits are annotation-only: no field is added, removed, reordered, retyped, or given a new default, so the FlatBuffers wire format and the generateddata_objects/*_generated.hheaders are unchanged.ASIC Coverage
ASIC-independent; passing PR CI is sufficient. No kernel selection, support surface, or default behavior changes — the matcher has no runtime consumer, and the new code is graph/schema logic with no device code and no arch-conditional paths. No multi-arch sweep required.
Testing Summary
$kernelmetadata namespace.= nulldefaults, scalar attribute types, and accessor round-trips reading live flatbuffer values.nlohmann::jsondata source backing the expression layer.umd_registry_genexit non-zero rather than emit a wrong binding.HIPDNN_ENABLE_KERNEL_INGESTOR=ON, so this change is verified alongside the ingestor foundation it will plug into, not just againstdevelop.Testing Checklist
ninja -C build && ./bin/hipdnn_plugin_sdk_tests- Status: Passed (197/197, including the ingestor suites from the parent branch)./bin/hipdnn_flatbuffers_sdk_tests- Status: Passed (187/187, including 8TestUmdOpSchemaRegistry)./bin/hip_kernel_provider_tests- Status: Passed (603 ran, 602 passed, 1 skipped forHIPDNN_ENABLE_SDPA=OFF; 121/121 acrossTestUmdCompiler,TestUniversalGraphMatcher,TestBindingContext,TestJsonLogic,TestJsonDataSource)./bin/hip_kernel_provider_integration_tests- ASICs: gfx1151 (Radeon 8060S) - Status: Passed (13 tests, process exit 0)cmake --preset hip-kernel-provider -G Ninja -DHIPDNN_ENABLE_KERNEL_INGESTOR=ON && ninja- Status: Passed (exit 0, no warnings introduced)graph.bfbs; working tree clean at the committed header - Status: Passedumd_registry_genontests/umd/fixtures/bad_nonlong.fbs- Status: Passed (rejects with a diagnostic, exit 1)pre-commit run --files <changed>- Status: PassedTechnical Changes
dnn-providers/hip-kernel-provider/include/hip_kernel_provider_common/umd/:UmdCompilervalidates ahipdnn.umd/v1descriptor and compiles its criteria to a reusableCompiledUmd;UniversalGraphMatcherruns oneCompiledUmdagainst a live flatbuffer graph, resolving pattern nodes by opcode and enforcing that each pattern variable binds to one tensor across all nodes;BindingContextis the resulting queryable symbol table and doubles as the JsonLogic data source.UmdPathParsekeeps the compile-time and runtime path resolvers from drifting.umd_registry_gen, a host tool that readsgraph.bfbsand emits the strongly-typed op-schema registry (RFC 0018 Appendix B). Annotations survive only in the reflection schema, soflatc --cppalone cannot produce this; CMake now emits.bfbs, builds the tool, and regeneratesumd/op_schema_registry_generated.hppinto the committed include tree underHIPDNN_GENERATE_SDK_HEADERS. The tool fails closed on an annotation violation instead of emitting a wrong binding.umd_opcode/umd_input_tensor/umd_output_tensor/umd_nameonce indata_types.fbsand apply them toSdpaAttributesand its UID fields. Metadata only; no field, ordering, type, or default changes, so the wire format and the generated data-object headers are unaffected.flatc-hipdnnpre-commit hook only refreshes the--cppdata-object headers (it cannot compile a C++ tool), so a schema change committed via pre-commit could leave the registry stale. The superbuild workflow now fails if the post-build tree is dirty at that header, and the split of responsibility is documented in the hook config andrun_flatc.py.JsonLogic'snlohmann::json-backed data source and linknlohmann_jsonexplicitly into the provider test suite so it also builds standalone and withHIPDNN_SKIP_JSON_LIB=ON.umd/*_generated.hppfrom clang-format and clang-tidy, scoped to that directory so the hand-writtenOpSchemaRegistry.hppbeside it stays checked and formatted.umd_binding_model,umd_matcher_pipeline,umd_symbol_binding), plus the PoC plan underprojects/hipdnn/docs/rfcs/.