Repository navigation
[hip-kernel-provider] Save and load execution plans of generic ingestor engines (ALMIOPEN-2642) - #13315
Merged
adickin-amd merged 44 commits intoOct 23, 2026
Conversation
github-actions
Bot
added
documentation
project: hipdnn
project: hip-kernel-provider
labels
⏳ Waiting on Required Checks⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
⏳ Waiting for required checks to start or finish:
|
Add the byte format that a saved ingestor execution plan becomes, and a
codec that encodes and decodes it. No plan, handler or GPU is involved yet;
later changes fill the payload at save time and consume it at load time.
The payload is a 48-byte header followed by a FlatBuffer body:
- header: marker, format major/minor, kind, header size, reserved word, and
a SHA-256 of the body. The 48-byte size keeps the FlatBuffer 8-aligned.
- the reader checks in a fixed order: header, then the whole-body digest,
then the FlatBuffers verifier, then structural checks, then compatibility.
Damage is reported as INVALID_VALUE before any NOT_APPLICABLE
("valid but not loadable here"); version and kind come first because they
define how the rest is read.
- version rule: same major, any minor up to the reader's own.
- launch inputs are named, typed values; the runtime pass-by-value UID list
is always written and a non-empty list is refused.
The codec is written for every kernel source kind. Saving kpack-only plans
is a scoping choice for this iteration enforced elsewhere, not a limitation
of the format.
The generated header is committed, produced by flatc 25.9.23 (the version
its static_assert requires). With HIPDNN_GENERATE_SDK_HEADERS on it is
regenerated using hipDNN's flatc resolution order; a forced regeneration
was verified byte-identical. Lint/format configs exclude only that file;
the .clang-tidy change lands in hipDNN's file because the providers
symlink to it.
Tests: 30 GPU-free cases pass on gfx90a and gfx942; each check test was
also run in a failing configuration and seen to fail. clang-tidy has not
yet been run over these files.
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…HEADERS option Add the provider's own option, HIPKERNELPROVIDER_GENERATE_SERIALIZE_HEADERS (default ON), to control regeneration of the committed ingestor plan FlatBuffers header (serialization/ingestor_plan_generated.h) from its schema. It behaves like hipDNN's HIPDNN_GENERATE_SDK_HEADERS switch, but applies only to hip-kernel-provider, so the provider header no longer depends on a hipDNN SDK option. The flatc resolution order and flags are unchanged. Wrap the flatc resolution and the regeneration rule in block()/endblock() so the HKP_* paths and _hkp_* temporaries go out of scope automatically, and drop the explicit unset() calls. The custom command, custom target and dependency are global and stay effective outside the block. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…hema Rename and tighten the payload schema before any plan is saved with it. The format version stays 1.0, because no payload exists in the old layout. - The root table IngestorPlan is now ExecutionPlan. Its namespace keeps it apart from hipDNN's SerializedExecutionPlan. - engine_id (long = null) is replaced by engine_name (string, required): the engine's readable name. A loader can recompute the id from the name with engineNameToId. The writer refuses an empty name (INTERNAL_ERROR). - kernel_id is renamed kernel_descriptor_id. - The fields launch and kernel are renamed dispatch_data and kernel_image. The table type LaunchInputs is renamed DispatchData. - provider_version is required; the writer refuses an empty version (INTERNAL_ERROR). - The reader refuses a non-empty runtime pass-by-value list with "which this provider does not support" rather than "cannot restore". - Digest hex parsing accepts only lowercase, which is what sha256Hex writes. - The text sanitiser for payload-sourced messages is exposed as ingestorPlanMessageText for other refusal sites. - "signature" fields and members are marked as the kernel's argument list. The generated header is regenerated with flatc 25.9.23. Tests: the 32 Phase 1 cases pass on gfx90a and gfx942, and on the Phase 1 code built without any later change; each new check test was run in a failing configuration and seen to fail; clang-tidy reports no diagnostics. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Add everything needed to turn a built ingestor plan into the saved-plan payload. Nothing is wired into hipDNN yet; that comes with the plugin hooks. - Plugin SDK seams, each defaulting to "not serializable": a handler save/restore pair (SavedDispatch.hpp, IKernelDispatchHandler), an IPlan saveable-plan method, GenericPlan accessors and BenchmarkPlan's chosen winner. A benchmarking plan is saved only once a winner exists. - captureIngestorPlan fills the payload: engineName, kernelDescriptorId, the plan's own workspace size, the dispatch data and the kernel image. The caller passes both the engine id and the engine name; a name whose engineNameToId hash is not the id is refused as INTERNAL_ERROR, and the message names the name and the id in hex (formatEngineIdHex). - Kpack bytes are re-read from the archive at save time and checked against the descriptor's SHA-256; the plan keeps the decorated device arch so the same archive entry is re-selected. Nothing extra is held in memory. - Each pack's launch inputs are named, typed values, declared in a per-engine header beside the pack (PointwiseLaunchValues.hpp, ConvFwdLaunchValues.hpp, Gfx950AttentionDenseLaunchValues.hpp), and each handler is also registered under a versioned alias (.v1); descriptors keep the unversioned name. - The plugin SDK gains toString(KernelSourceKind) in Descriptors.hpp, the inverse of kernelSourceKindFromString; capture uses it for its refusal text. Two SDK tests cover the round trip and the "unknown (n)" fallback. - One source-kind gate decides what can be saved. Only kpack kernels pass today because embedded_source kernels have no recorded argument signature and their bytes are not exposed. The refusal says the source kind is one "which saving does not yet support": a scoping choice for this iteration, not a limitation of the format; hipRTC support is expected. The SDK header changes need hipDNN-core codeowner review. Tests: the 21 GPU-free and 8 GPU provider cases, the 4 SDK cases and the 32 Phase 1 cases pass on gfx90a and gfx942, on the Phase 2 code built without any later change; the two source-kind SDK tests and the capture cases pass on gfx942. The quick set (without the external-integration suites) has no new failures. Each check test was also run in a failing configuration. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Add restoreIngestorPlan, the provider function that the plugin deserialize hook will call. It rebuilds an executable plan from the payload bytes alone: no graph, no matching, no descriptor content other than the engine names, no archive, no benchmark result and no compile. Nothing is wired into hipDNN yet; that comes with the plugin hooks. - Restore runs nine checks in a fixed order. The first failure refuses the payload, and each message names its cause and the kernel: header, body digest, verifier and field structure, engine identity, device and code-object target, dispatch handler, kernel argument signature, launch-value contract, and module load. Damage gives INVALID_VALUE. A valid payload that cannot run here gives NOT_APPLICABLE. A handle with no device gives INTERNAL_ERROR. - Engine identity: the saved engine name hashes to an engine id, and loadedIngestorEngineName must find a loaded ingestor engine with that id and the same name. - Target rule (CodeObjectTarget.hpp): a kpack target passes on each device that could select its archive entry. Every other source kind needs the exact target, feature flags included. - loadCodeObjectOnDevice (compilation/CodeObjectLoad) loads bytes as a module on one device and resolves the symbol, with no module cache. A failure names its stage and keeps the HIP error. The provider does not clear HIP's last error after a refused load. - IngestorKernelCode::fromSavedKernelCode loads the saved code on the handle's device during restore, so a driver refusal fails the restore and not a later execute. Another device loads its own module from the same bytes on first use, after an exact-target check. The bytes stay in memory for the life of the plan. - IngestorKernelCode refuses an entry with no program or no kernel (INTERNAL_ERROR). Every write to its per-device memo goes through that check. - Each pack gets a restoreLaunch override and a reader for its launch values, in its own launch-value header. The reader accepts only the versioned dispatch name, and refuses a missing, extra, wrong-type or out-of-range value. prepare() and restore build the dispatch through the same makePrepared function, so both give the same geometry. - RestoredIngestorPlan reports the payload's workspace size and refuses a null workspace when it needs one. Capture refuses to save a restored plan and tells the caller to keep the original bytes. - Errors: restore and capture throw HipdnnPluginException with a precise status. A code-object load failure becomes a refusal inside IngestorKernelCode. No other exception is caught in this code; the plugin boundary converts each exception to a status. - The plugin SDK gains toString(MetadataType) in Descriptors.hpp, the inverse of metadataTypeFromString; restore messages use it for value types. SavedKernelCode gains kernelId, so messages name the kernel. New SDK tests cover the MetadataType round trip, the "unknown (n)" fallback, the IPlan saveablePlan default, and a benchmarking plan whose chosen candidate has no saveable plan. The SDK header changes need hipDNN-core codeowner review. Tests: - GPU-free cases cover each restore check, the check order for each adjacent pair, the target rule, each pack's reader in both directions, the restored plan, the refusal to re-save, and the engine lookup, including a loaded engine that no descriptor set defines. - GPU-free reload cases use a test subclass of IngestorKernelCode in place of the driver load. They cover a load on a new device, reuse of a known device, refusal of a device that cannot run the exact target, and refusal of an entry with no kernel. No test needs more than one GPU. - The round trip captures, encodes, restores and executes a kpack Pointwise plan and a kpack ConvFwd plan. Each checks the original output against exact expected values, then checks that the restored plan, run on a new handle after the original plan is destroyed, gives the same bytes. A second case restores with the archive deleted and the module caches cleared. Other cases refuse code bytes the driver rejects and a symbol the code object lacks. - The packed-plan tests share one helper header and one stub dispatch handler, and describe each packed set with a case object in its own header. They use the test SDK tensors and workspace. - The test helpers read descriptors with the plugin SDK kernel parser. The refusal helper fails a test that throws any other exception type and names the type and the message. - The unit binary reads its own packed copy of archive_fixture, staged outside its discovery root. The integration suite corrupts only its own copy, so the two binaries need no resource lock. Tests: the quick set (without the external-integration, asm-sdpa external-integration and gpu_ref suites) passes 36/36 on gfx90a and gfx942. The targeted provider cases (148) and plugin SDK cases (19) pass on gfx942. Each new or changed check test was also run in a failing configuration and failed as expected. clang-tidy reports no diagnostics. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
hip-kernel-provider now exports the three plugin hooks that hipDNN uses
to save and load an execution plan:
hipdnnEnginePluginSerializeExecutionContext,
hipdnnEnginePluginDestroySerializedExecutionContext and
hipdnnEnginePluginCreateExecutionContextFromSerialized. They are compiled
only when the kernel ingestor is enabled. Each runs inside the plugin's
exception boundary, so a refusal returns a status and a message and no
exception crosses the C interface. hipDNN's plugin API is unchanged.
The save hook receives no engine ID from hipDNN, so the provider records
it. Container wraps each ingestor engine in a forwarding engine that,
after building an execution context, stores the engine config's engine
ID on the provider Context. The load hook stores the engine ID of the
restored payload; restoreIngestorPlan now returns it with the plan.
Contexts of the provider's other engines carry no engine ID.
Save reads the engine ID from the context and the engine name from
hipDNN's engine-name registry (tryGetEngineNameFromId), which the
descriptor loader fills for every ingestor engine. A context with no
engine ID, or an ID with no registered name, is refused with
NOT_APPLICABLE ("does not support saving execution plans"). Otherwise the
plan is captured and encoded as the saved ingestor plan payload; the
existing capture refusals (a benchmarking plan with no winner yet, a
restored plan, a kernel kind that cannot be saved) apply unchanged. Load
restores the plan and wraps it in a new execution context.
The hipDNN data SDK gains tryGetEngineNameFromId, which returns the
registered name for an engine ID or no value, and never throws
(getEngineNameFromId throws std::out_of_range for an unregistered ID).
engineNameOrHex is unchanged. The descriptor loader's comment on
engine-name registration now states only why a throw there is safe to
ignore.
Restore's engine check now takes the engine name from the registry
instead of scanning the discovered descriptor sets: the payload's engine
name must hash to an engine ID that the handle has loaded, and the
registry must map that ID back to the same name. loadedIngestorEngineName
is removed.
Each ingestor engine now reports SUPPORTS_EXECUTION_PLAN_SERIALIZATION
when every kernel of every pack in its descriptor set passes the
source-kind gate. Container passes the computed descriptor set to the
engine; the discovered sets are unchanged. From this commit, to_binary()
of a built plan from a noted engine embeds the plan.
The provider's other engines get no note. Their plan-only save now fails
with HIPDNN_STATUS_PLUGIN_ERROR (NOT_APPLICABLE) instead of
HIPDNN_STATUS_NOT_SUPPORTED, because the hooks are per plugin.
The packed-plan test helpers runPlan, stubKpackKernel and makeStubPlan
move into the shared PackedPlanTestSupport.hpp.
Tests: TestPluginSerialization (hook null checks, save and load
refusals including a context with no engine ID, and a ConvFwd round trip
through the hooks that checks exact values, fills the output with a
sentinel before each restored execute, and checks the restored context's
engine ID and its re-save refusal); TestKernelIngestorEngine (the
forwarding engine, which contexts carry an engine ID, the gate, the note,
and a guard that every dispatch handler of a noted engine saves and
restores its launch inputs); TestIngestorPlanRestoreChecks (the registry
engine check); TestContainer.EngineNameRegistryHoldsEveryIngestorEngineName;
TestEngineNames.TryGetEngineNameFromId* (data SDK).
Testing: on gfx942, 29 of 29 failing configurations are caught by the
tests above (the clean tree passes 72 of 72, before and after). The
quick test set, excluding the three standing suites, passes 37 of 37
on gfx90a and on gfx942. clang-tidy on a fresh gfx90a tree reports no
diagnostics. The gfx950 build passes; the gfx950 test job is still
queued.
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The provider keeps no record of which engine built or names a saved
plan. hipDNN's envelope engine ID is the only engine check at load:
hipDNN routes the load to the plugin that serves that ID, and refuses an
ID that no loaded plugin serves.
Removed: the engine_name field of the saved plan payload and
IngestorPlanPayload::engineName; the engine ID and engine name
parameters of captureIngestorPlan and its check that the name hashes to
the ID; restore's engine check, its loadedEngineName lookup and
IngestorPlanRestoreResult (both restoreIngestorPlan overloads return the
plan, and IngestorPlanRestoreEnvironment holds only the device
resolver); Context::engineId and setEngineId; the forwarding engine
detail::EngineIdRecordingEngine and its wrapping in Container; and the
save hook's engine ID and engine-name registry lookup. The load hook
stores only the restored plan. The data SDK helper
tryGetEngineNameFromId and its tests stay.
Save now classifies a plan by its type. Capture refuses a restored plan
first. A plan with no saveable single-kernel plan gets the "has not
chosen a kernel yet" refusal when it is a BenchmarkPlan, and otherwise
the non-ingestor refusal ("this plan was not built by an ingestor
engine; its engine does not support saving execution plans").
refuseNonIngestorPlanSave owns that text. The save hook refuses a
context without a plan with the same refusal, before it reads the plan.
The source-kind gate and the handler save follow, unchanged.
Restore runs eight checks in this order: header, body SHA-256,
FlatBuffers verifier and field structure, device target, dispatch
handler, kernel signature, launch values, module load. A refusal for an
unknown dispatch name names the saved kernel.
The payload format version stays 1.0, and the schema drops engine_name.
ingestor_plan_generated.h is regenerated from the schema with flatc
25.9.23. A payload saved before this commit has the same version but a
different table layout, so the FlatBuffers verifier refuses it as
damaged (INVALID_VALUE, "ingestor plan data is damaged").
Known limits: an engine that is not installed fails at load with
hipDNN's INTERNAL_ERROR "Invalid engine ID" instead of a provider
message; an engine that is advertised but failed to construct still
loads its saved plans, because restore does not use the engine; nothing
checks that a payload belongs to the envelope's engine; and a payload
does not name its engine.
The <optional> includes in KernelIngestorEngine.{hpp,cpp}, which only
the removed engine lookup used, are gone.
Tests: new TestIngestorPlanCapture.RefusesAPlanNoIngestorEngineBuilt,
TestPluginSerialization.SaveRefusesAPlanNoIngestorEngineBuilt and
SaveRefusesAContextWithoutAPlan, through a shared NonIngestorPlan stub
in PackedPlanTestSupport.hpp, which also holds the two refusal phrases.
RefusesABenchmarkPlanWithoutAWinner checks that the non-ingestor text is
absent; SaveLeavesTheOutputUntouchedOnRefusal saves a non-ingestor plan
in its second half; TestIngestorPlanRestoreChecks.RefusesAnUnknownAlias
checks that the refusal names the saved kernel ID. The round-trip tests
restore through the production overload. The codec tests are renamed to
VerifierRejectsABodyWithoutAProviderVersion and
WriterRefusesAnEmptyProviderVersion. Removed: the engine-name capture
test, the three restore engine tests and the EngineBeforeTarget order
case, the six engine-ID tests in TestKernelIngestorEngine, and
TestContainer.EngineNameRegistryHoldsEveryIngestorEngineName.
Testing: on gfx942, 32 of 32 failing configurations are caught by the
targeted tests (the clean tree passes 104 of 104, before and after).
The quick test set, excluding the three standing suites, passes 37 of
37 on gfx942. clang-tidy on a fresh gfx90a tree reports no diagnostics.
The gfx90a and gfx950 builds pass; their test jobs are still queued
because no node was free.
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Tests save and load through the public hipDNN API, in one process. Each
round trip saves with one hipDNN handle, destroys the original graph and
that handle, clears the provider's code caches, creates a new handle
(asserting that the provider creates a new Container), loads, and runs
the loaded plan twice against the CPU reference taken from the original.
The load is checked to run no kernel selection or benchmarking, against
a positive control on the original build. The loaded plan reports the
engine ID and workspace size of the original.
- IntegrationGpuKernelIngestorSaveLoad: hipkernel:Pointwise and
hipkernel:ConvFwd, each saved as graph and plan
(to_binary/from_binary) and as plan only
(to_compiled_plan_binary/from_compiled_plan_binary).
- IntegrationGpuGfx950AttentionDenseSaveLoad:
hipkernel:Gfx950AttentionDense in both forms; it runs on gfx950 only.
- IntegrationGpuKernelIngestorSaveLoadRefusals: a benchmarking plan
before its first execute ("has not chosen a kernel yet"), a plan that
was itself loaded ("loaded from a saved payload", for both forms), a
payload with one code-object byte changed ("body digest mismatch"),
and a payload with another format major version. Each case checks the
frontend error code, the plugin status and the cause phrase.
The provider's graph verification harness is split into
prepareVerification and executeAndValidate, so a test can validate a
loaded graph against the original's reference. The GPU run comes before
the CPU reference, so to_binary() never runs before a plan's first
execute. When a test reuses a verification run, executeAndValidate
fills the GPU outputs with a sentinel before it executes, so a loaded
plan that writes nothing cannot pass on the original run's values.
Single-run harness callers are unchanged. Validators are shared, so a
caller's comparator survives repeated calls. The module-cache reset
helper moves to the new IngestorIntegrationSupport.hpp.
Documentation:
- RFC 0009 gets a corrected Plugin API block with the three hooks as
EnginePluginApi.h declares them, including the three-argument
hipdnnEnginePluginSerializeExecutionContext, and states that hipDNN
uses the hooks only when a plugin exports all three.
- RFC 0009 gets a hip-kernel-provider ingestor section: the supported
engines, the payload layout (format version 1.0), versioning, the
load checks in order (including that the dispatch handler supports
restore), the statuses and message prefixes, the save rules in order,
and security. hipDNN routes a load by the envelope engine ID; nothing
checks that the payload belongs to that engine.
- The provider README gets "Saving and Loading Execution Plans".
- The plugin development guide lists the three hooks and describes
saving from ingestor engines.
Testing: on gfx942, the save/load tests pass (the two gfx950-only cases
skip), as do the capture unit tests and the harness integration tests;
9 of 9 failing configurations that can run on gfx942 are caught. The
quick test set, excluding the three standing suites, passes 37 of 37
on gfx90a and on gfx942. clang-tidy on a fresh gfx90a tree reports no
diagnostics. The gfx950 save/load run and its two failing
configurations are still queued because no gfx950 node was free.
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The saved-plan tests check each rule once, in the test that sees it end
to end. No product code changes.
Unit tests:
- TestPluginSerializationKpack.RoundTripsAKpackPlanThroughTheHooks runs
for the Pointwise and the ConvFwd packed sets. One plan goes through
its whole life through the C API hooks, one SCOPED_TRACE per stage:
the original run gives the exact expected output; the save hook
returns bytes; once the archive is removed, a second save is refused
with OPEN_ARCHIVE and the archive path, and the output is untouched;
the original plan still computes the same bytes; the original is
destroyed and the module caches are cleared, then the load hook
restores the plan on a new handle, without the archive, with the
original workspace size; two restored runs give the original bytes
and the exact output, and load nothing into the module caches; a
re-save of the restored plan is refused ("loaded from a saved
payload") with the output untouched; the destroy hook frees the
bytes.
- At the hook level, the GPU-free tests are the null-argument checks,
the refusal of a plan that no ingestor engine built and of a context
without a plan (both leave the output untouched), and the refusal of
damaged bytes. The benchmarking-plan and restored-plan refusals are
tested where they are made, in captureIngestorPlan.
- TestIngestorPlanCaptureGate.AppliesTheSourceKindGate needs no device.
It uses stub plans whose stub kernel has the source kind of the
parameter: an embedded_source kernel is refused at the gate, and a
kpack kernel passes the gate and is refused at the next step, by the
stub dispatch handler. stubKernel(seed, kind) makes the stub kernel;
makeStubPlan takes the kind, with kpack as the default.
- TestIngestorPlanRoundTrip keeps the two restore refusals for code
bytes or a symbol that the driver does not accept.
Integration tests:
- Each IntegrationGpuKernelIngestorSaveLoad round trip also checks that
the loaded plan, in its new handle, cannot be saved again
("loaded from a saved payload"), in both save forms and for both
engines.
- IntegrationGpuKernelIngestorSaveLoadRefusals:
RefusesToLoadDamagedBytesOrAnotherFormatMajorVersion saves one
ConvFwd plan, then loads two changed copies, each under its own
SCOPED_TRACE: one with a code-object byte changed ("body digest
mismatch") and one with another format major version.
Removed tests (the tests above make their checks, except the six
checks in the next paragraph):
- TestIngestorPlanRoundTripKpack (both packed sets):
RestoredPlanComputesWhatTheOriginalComputes and
RestoresWithoutTheArchive.
- TestIngestorPlanCapture: RefusesWhenTheArchiveIsGone.
- TestPluginSerialization: SaveRefusesABenchmarkPlanWithoutAWinner,
SaveRefusesARestoredPlan and SaveLeavesTheOutputUntouchedOnRefusal.
- IntegrationGpuKernelIngestorSaveLoadRefusals:
RefusesToSaveAPlanItLoaded, RefusesToLoadDamagedBytes and
RefusesToLoadAnotherFormatMajorVersion.
TestPluginSerialization.RoundTripsAKpackPlanThroughTheHooks becomes the
parameterised TestPluginSerializationKpack test.
The checks with no remaining test are a round trip through
restoreIngestorPlan without a Container, a restore while the original
plan is alive, the archive refusal at the capture level on a GPU, the
benchmarking refusal text at the hook level in the unit tests, the gate
check with a real packed kpack plan (CapturesAKpackPlan still captures
one), and the refusal of a real embedded_source Pointwise plan.
Test counts: the three unit files go from 25 to 18 tests
(TestPluginSerialization* 8 to 6, TestIngestorPlanRoundTrip* 6 to 2,
TestIngestorPlanCapture* 11 to 10); the integration *SaveLoad* tests go
from 10 to 8, of which 2 run on gfx950 only.
Testing: the targeted unit tests pass 66 of 66 on gfx942 and on gfx950;
the save/load integration tests pass on gfx942 (the two gfx950-only
cases skip) and 8 of 8 on gfx950. All 27 failing configurations are
caught: 25 on gfx942 and the two gfx950-only ones on gfx950. The quick
test set, excluding the three standing suites, passes 37 of 37 on
gfx942; the gfx90a test job is still queued (its build passed).
clang-tidy on a fresh gfx90a tree reports no diagnostics.
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
A saved gfx950 attention_dense plan is held to the same softmax scale range as plan build. readGfx950AttentionDenseLaunchInputs first checks that the saved scale is a finite double that a float holds exactly, then refuses a scale outside [2^-64, 2^4] as incompatible (NOT_APPLICABLE). Plan build and restore share one predicate, isSupportedAttentionScale. The plan-build decline and its log text are unchanged. TestPackLaunchValues checks both bounds and four scales outside them without a GPU. ReportsTheSerializationNoteOnlyOnSaveableEngines checks both ways a Container builds an ingestor engine. After a rediscovery, the first Container takes the state manager that discovery built, and the test asserts that the slot is empty; a later Container builds its own. Each must report the serialization note once on ConvFwd and never on Pointwise. Testing: on gfx942 and gfx950 the targeted unit tests pass 113 of 113, the engine and container tests pass 20 shuffled repeats, and the save/load integration tests pass (8 of 8 on gfx950). 41 of 41 failing configurations are caught, the two gfx950-only ones on gfx950. The quick test set, excluding the three standing suites, passes 39 of 39 on gfx90a and on gfx942. Builds with serialize-header generation off and with rocKE off succeed. clang-tidy on a fresh gfx90a tree reports no diagnostics. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
KpackModuleCache.hpp has one kpack code-object reader, readVerifiedKpackCodeObject(archivePath, tocKey, deviceArch, expectedSha256, selectedArch). It opens the archive through SharedKpackArchives::open, selects the arch, reads the entry and checks its digest. It evicts the kept handle before it throws a decode failure or a digest mismatch, and keeps it on a missing arch or entry. This is the eviction rule that KpackModuleCache::load applied on develop. KpackModuleCache::load() and the save path (IngestorKernelCode::readCodeObject) both call this reader; load() calls it without a try/catch. The reader overload that took an opened archive is removed, and SharedKpackArchives::openUnretained is private again; the SharedKpackArchives class text is the same as on develop. While a lease is held, which a hook save always has, a save reads the archive through the handle that the load kept. So a save after the archive was deleted can succeed from the kept handle; the digest still checks the bytes. Nothing requires a save to fail when the archive is removed after the plan build: - The hook round trip loses its "save after the archive is removed" and "original run after the archive is removed" stages. It removes the archive in "load without the archive", after the plan is destroyed and the module caches and kept archives are reset, so the restore still runs without the archive. - No test checks a save refused because the archive was deleted after the plan build, or the C API status of a save refused for an unreadable archive. RefusesToSaveWhenTheArchiveIsAbsent (an archive that never existed) still checks the refusal of a failed read. The README rule and RFC 0009 say that an unreadable archive or changed bytes fail the save, and that the save hook reads the archive at the path from the plan build. The readCodeObject comment says that it takes no lock of the kernel, and the eviction comment in TestKpackModuleCache.cpp points at readVerifiedKpackCodeObject. Testing: on gfx942 the targeted unit tests pass 126 of 126 and the save/load and kpack integration tests pass (the two gfx950-only cases skip). 20 of 20 failing configurations are caught. The quick test set, excluding the three standing suites, passes 39 of 39 on gfx942; the gfx90a test job is still queued (its build passed). clang-tidy on a fresh gfx90a tree reports no diagnostics. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
CMiservaAMD
force-pushed
the
users/cmiserva/almiopen-2642/save-execution-plans
branch
from
October 9, 2026 16:25
3fd2ef6 to
6a7e858
Compare
The serialize hook released ownership of the payload into the caller's struct before its final log call. That log builds a string and can throw, which tryCatch turns into a failure status; hipDNN then never calls the destroy hook and the payload leaks. Logging first makes the release the last statement, so no fallible call follows the hand-over. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The probe scored any std::exception as evidence that a handler overrides save and restore, so an override that threw for an unrelated reason passed. Restore now asserts the shared restore checks' NOT_APPLICABLE refusal. Save catches std::bad_cast alone, which is what a correct handler raises on a foreign prepared dispatch: its overrides open with the down-cast and carry no refusal path, so no status distinguishes them. Any other exception type now propagates and fails the test. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Without an explicit id a field's vtable slot comes from its position, so inserting a field renumbers every later one while the Verifier still passes and values silently swap. With ids the slot comes from the number, and a new field takes the next unused one wherever it is written. The ids reproduce the current numbering exactly: regenerating ingestor_plan_generated.h from this schema yields a byte-identical file. The union in LaunchValue occupies two slots, so its declared id names the value and the slot below holds the implicit type byte. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…egenerated Generation was gated on the option, a resolved flatc and the flags file in one condition, and any of them failing skipped the step in silence. A schema edit could then ship against the committed header with nothing to show for it. The option being off now reports that the committed header is used as-is. With the option on, an unresolved flatc or a missing flags file is a configure error naming which one failed, matching how hipDNN's own FlatBuffersGenerate.cmake treats a missing flatc. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
getEngineNameFromId and engineNameOrHex each repeated the registry lookup, and engineNameOrHex reached its hex fallback by catching out_of_range, so the common logging path ran on exceptions. Both now call tryGetEngineNameFromId, which until now had no caller outside its own test. The thrown message is unchanged. engineNameOrHex had no test despite five callers; its two branches and the hex in the thrown message are now covered. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The digest is checked before the verifier, so any corruption that is not resealed stops at the digest and the structural layer below it was reached only by three hand-written cases. The sweep flips each body byte, reseals, and requires the decode to refuse with a damaged or incompatible status whose message carries the matching category prefix, or to return a payload whose lengths all fit the body it came from. Any other exception type fails. The hand-written cases stay: a sweep cannot produce a missing required field, which needs a rebuilt vtable rather than an altered byte. The out-of-bounds reads this can expose are only caught under a sanitizer build; elsewhere it checks status, category and structure alone. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
… samples Each pack reads its launch inputs from a versioned set of named values, and nothing held that set still: renaming, retyping, adding or dropping a value would change what an already-saved plan means with no test to notice. Each sample is bytes an earlier build wrote. The test decodes them and asserts the launch inputs, so it reads in one direction only; comparing freshly encoded bytes would fail on a flatc or builder layout change that means nothing. The gfx950 sample also asserts the geometry its block size derives, which catches a value that keeps its name and type while changing what it denotes. Pointwise and conv build their prepared dispatch inside an anonymous namespace, so their geometry is out of a test's reach and only the wire contract is pinned. A failing sample is never fixed by regenerating it. Regenerated bytes assert that today's code reads today's output, which is not compatibility at all. When the format changes, add a sample beside the old one and leave the old one passing; a sample that can no longer be read is the finding. The disabled printer produces new samples, not replacements for existing ones. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
… comments The launch-contract samples were written by the build that introduced contract v1, not by an earlier build; the file header now says so. The three leftover paste markers are removed, and two comments now state only why the dispatch and value names are spelled out. The readSample comment no longer claims a partly decoded sample skips the value assertions, since EXPECT_EQ does not stop the test. The regeneration comment in the ingestor CMakeLists.txt now states the rule alone: with the option ON, a missing flatc or flags file is a configure error. Comment-only change. Testing: on gfx942 (MI300X), after a full build of debug-gfx942, the quick test set excluding the three standing suites passes 39 of 39 (hip_kernel_provider_tests_quick_suite 7.38 s), and the launch-contract, codec, engine and plugin-serialization tests pass 48 of 48 with none skipped. A configure with the option ON and no flatc stops at the new configure error; with the option OFF it reports the committed header and continues. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
CMiservaAMD
marked this pull request as ready for review
October 13, 2026 15:34
Conflict resolutions (all in dnn-providers/hip-kernel-provider/): - CMakeLists.txt: keep both options, HIPKERNELPROVIDER_REQUIRE_DVC_DESCRIPTORS and HIPKERNELPROVIDER_GENERATE_SERIALIZE_HEADERS with its TheRock default. - src/compilation/KpackModuleCache.hpp: load() keeps the single reader, readVerifiedKpackCodeObject. The makeKey comment names that reader. - src/engines/kernel_ingestor_engine/CMakeLists.txt: keep Gfx1151WmmaAttentionNative.cpp and the two serialization sources. - src/tests/CMakeLists.txt: keep TestKpackGenericArchive.cpp and TestKpackCodeObjectRead.cpp. - src/tests/engines/kernel_ingestor_engine/CMakeLists.txt: keep TestGfx1151WmmaAttentionMatchers.cpp and the saved-plan tests. selectKpackArchiveArch ranks archive entries by entryTier (an explicit key beats a generic one; ties keep archive order). KpackModuleCache::selectArch forwards to it, so load and save pick the same entry. TestKpackArchSelection needs no change: for non-generic keys the tier rule equals the prefix rule. Known red on gfx11-member builds only: the packed gfx1151 WMMA attention set gets the serialization note but its handler cannot save, so TestKernelIngestorEngine.EveryHandlerOfANotedEngineSavesAndRestores fails there. The next commit that gates the note on the handler fixes it. Not visible on gfx942. Testing (gfx942 MI300X, node ctr-cx63-mi300x-12, image hipdnn_20261017_gfx942): - Baseline, clean dev/develop 299f298: build OK, quick set 53/53 (job 68125717). - Merged tree: full build OK, quick set 53/53 (job 68125862). - Targeted provider tests: unit 187/187, IntegrationGpuKernelIngestor* 26/26 (job 6812586). - Plugin SDK targeted tests: 231/231 (job 68125865). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…for a generic archive entry
A saved kpack plan recorded the archive key that served the device as its
target. The packer writes every archive key as a bare name with no feature
suffix. For an explicit entry, the key therefore equals the device base name.
For a generic entry (for example gfx11-generic serving gfx1151), the key is
the family name, and neither restore check admits a family name: the load-time
target check and the per-launch architecture check both compare against a
base architecture.
The capture now records the base name of the device the plan was saved on
(stripArchFeatures of the device arch). For explicit entries the value is
unchanged. For a generic entry it pins the plan to the saving GPU type, which
both restore checks then admit unchanged. A plan does not move to another
member of the generic family.
No format change: the schema carries no comment on target, so
ingestor_plan.fbs, the generated header and the schema hash are unchanged.
The doc comment of SavedKernelCode::target in the public plugin SDK header
(SavedDispatch.hpp) changes, so it needs hipDNN plugin SDK codeowner review.
The docs of IngestorCodeObject::target, IngestorPlanPayload::target and
isCodeObjectTargetCompatible change to match.
The restore refusal for an incompatible device now says the plan was saved
for the target, not that its code object was built for it.
Tests:
- TestKpackArchSelection: an explicit entry beats a generic entry; a generic
entry serves a member device.
- TestKpackCodeObjectRead.ACaptureRecordsTheDeviceBaseName (host-only): a
read from a copy of the gfx942 unit shard for device
gfx942:sramecc+:xnack- records gfx942.
- TestKpackCodeObjectRead.ReturnsVerifiedBytesWithoutLoadingAModule now
expects the stripped device name.
- TestIngestorPlanRestoreChecks.AdmitsAPlanSavedOnTheSameBaseArch and
TestIngestorKernelCodeReload.LaunchesARestoredPlanOnItsSavedBaseArchOnly.
Known gaps: the generic capture path is proven by review only. A gfx942 tree
holds no generic archive, so a test cannot tell the selected key from the
stripped device name there (recording the selected archive key again is not
caught on gfx942). There
is no gfx11 device round trip (no hardware).
Testing (gfx942 MI300X, image hipdnn_20261017_gfx942):
- Targeted unit tests 113/113 and IntegrationGpuKernelIngestor* 26/26
(jobs 68126121 and, after the message rewording, 68126218).
- ingestor_plan.fbs and ingestor_plan_generated.h unchanged (git diff
--exit-code).
- Mutations:
- Recording the undecorated device arch fails
TestKpackCodeObjectRead.ACaptureRecordsTheDeviceBaseName and
ReturnsVerifiedBytesWithoutLoadingAModule.
- Letting a later equal-tier entry win fails
TestKpackArchSelection.TheFirstMatchingEntryWins.
- Ignoring the tiers fails AnExplicitEntryBeatsAGenericEntry and
TestKpackModuleCacheSelectArch.PrefersTheExactDeviceArchOverAGenericKey.
- Recording the selected archive key again is not caught on gfx942.
- After restore: 113/113 and 26/26 (job 68126185).
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ion only when the dispatch handler can save
The provider added HIPDNN_BEHAVIOR_NOTE_SUPPORTS_EXECUTION_PLAN_SERIALIZATION
to an engine when every kernel of the engine was a kpack code object. The
kernel source kind alone does not show that a plan of the engine can be
saved. The gfx1151 WMMA attention engine is kpack on gfx11 builds (the packer
rewrites its rocke kernels to kpack), but its dispatch handler keeps the SDK
defaults: it does not save launch inputs and does not restore a launch. On a
gfx11 build that engine reported the note, so the frontend's to_binary() would
fail for its plans. The same mismatch made the existing handler test fail on
gfx11 builds after the develop merge; this commit clears that failure.
Plugin SDK: IKernelDispatchHandler gets a defaulted virtual
supportsSaving() const that returns false. A handler that overrides both
saveLaunchInputs and restoreLaunch must override supportsSaving() to return
true. ABI: IKernelDispatchHandler
<THandle> is a header-only class template that each plugin instantiates, and
hipDNN core does not call its virtuals across a binary boundary. The sibling
saveLaunchInputs and restoreLaunch virtuals are new in this unreleased change.
A defaulted virtual is source-compatible for out-of-tree handlers. This
touches a public plugin SDK header, so it needs hipDNN plugin SDK codeowner
review. PluginDevelopment.md says to override supportsSaving and that the note
needs it.
Provider: the Pointwise, ConvFwd and Gfx950AttentionDense dispatch handlers
return true. The gfx1151 WMMA attention handler is unchanged and reports
false. supportsExecutionPlanSerialization() keeps the source-kind check and
also requires that each pack's dispatch is in the set, that its symbol
resolves in the dispatch registry, and that the handler reports
supportsSaving(). It registers the native symbols first (once per process),
so it also works on sets that did not come from discovery. The engine header
doc and the provider README state the new rule.
Tests:
- TestKernelIngestorEngine.setWithKernels gives each pack a dispatch that
names hipkernel.conv_fwd.dispatch, so the existing gate and note tests keep
their meaning.
- NoNoteWhenAHandlerCannotSave: a test-registered stub handler (alone, or
beside a handler that saves) and the real gfx1151 WMMA attention handler
give no note.
- NoNoteWithoutAResolvableDispatch: a pack whose dispatch is missing from the
set, and a dispatch symbol that is not registered, give no note.
- EveryHandlerOfANotedEngineSavesAndRestores becomes the two-way
EveryHandlerReportsSupportsSavingExactlyWhenItSavesAndRestores. It probes
every handler reachable from the discovered sets and from a fixed list of
the four pack dispatch symbols (the registry has no enumeration). For each
handler, supportsSaving() must equal "saves and restores", and saving must
equal restoring. The fixed list must have one entry per ingestorPacks()
entry, so a new pack fails the test until its symbol is added. At least one
handler on each side must be probed.
- TestIngestorGenericPlan.HandlerSaveAndRestoreDefaultToNotSerializable also
checks the supportsSaving() default.
Follow-up candidate: save support for the gfx1151 WMMA attention engine (a
versioned launch contract, goldens, alias test) is a separate ticket.
RFC 0009 "Supported Engines" states the handler rule too.
Testing (gfx942 MI300X, image hipdnn_20261017_gfx942):
- Targeted provider tests: unit 51/51, IntegrationGpuKernelIngestor* 26/26
(job 68126252). Plugin SDK targeted tests 33/33 (job 68126253).
- ingestor_plan.fbs and ingestor_plan_generated.h unchanged.
- Mutations, each caught:
- Conv handler reports no saving: the two-way handler test, the
note tests and two conv save/load integration tests fail.
- Gate ignores supportsSaving(): NoNoteWhenAHandlerCannotSave fails.
- Gate accepts a missing or unregistered dispatch:
NoNoteWithoutAResolvableDispatch fails.
- SDK default returns true:
HandlerSaveAndRestoreDefaultToNotSerializable and the two-way test
fail.
- Gfx950 attention handler loses its override: the two-way test fails.
- After restore: unit 51/51, integration 26/26, SDK 33/33
(jobs 68126348, 68126349).
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Conflict resolutions (include lists only; both sides kept): - projects/hipdnn/backend/src/descriptors/ExecutionPlanDescriptor.hpp: <mutex> and <optional>. - projects/hipdnn/backend/src/descriptors/ExecutionPlanDescriptor.cpp: hipdnn_flatbuffers_sdk/utilities/Uuid.hpp and <mutex>. - projects/hipdnn/backend/tests/descriptors/TestExecutionPlanDescriptor.cpp: <atomic> <chrono> <condition_variable> <cstddef> <cstdint> <memory> <mutex>. The per-plan execution counter is an atomic changed only by recordExecution(). The serialized-plan cache mutex guards only that cache. Neither touches the other's state. Testing (gfx942 MI300X, image hipdnn_20261017_gfx942): - Full build OK; quick set 55/55 (53 as before plus 2 new rocKE tests from develop) (job 68131261). - Targeted provider tests: unit 194/194, IntegrationGpuKernelIngestor* 26/26 (job 68131263). Plugin SDK targeted tests 231/231 (job 68131264). - TestExecutionPlanDescriptor.* 39/39 (job 68131265). - ThreadSanitizer: TestExecutionPlanDescriptor.* 39 tests x20, 0 reports (build/tsan-hipdnn, job 68131295). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
CMiservaAMD
enabled auto-merge (squash)
October 21, 2026 01:02
Conflicts resolved, keeping the behaviour of both sides: - Container.cpp (#13028, UHD runtime): engine creation now takes the whole descriptor set. Pass withComputedBehaviorNotes(set) to it, so the plan serialization note and the prediction and selector inputs both survive. - ConvNative.cpp, PointwiseNative.cpp (#13028): include lists; keep both. - kernel_ingestor_engine tests CMakeLists.txt (#14158, multi-launch rocKE kernels): keep TestLaunchKernelDefinition.cpp and TestRestoredIngestorPlan.cpp. - TestKernelIngestorEngine.cpp (#13028): DispatchRegistry moved to NativeHooks.hpp; keep all tests of both sides. - IKernelDispatchHandler.hpp (#14158): include list; launchNames() and the save, restore and supportsSaving() hooks all stay. - TestBenchmarkPlan.cpp (#13028): include list and appended tests; keep both. - hipdnn-superbuild-ci.yml (#13385, tooling tests in CI): both lanes pass HIPDNN_ENABLE_TOOLING_TESTS and HIPKERNELPROVIDER_GENERATE_SERIALIZE_HEADERS. - PluginDevelopment.md, RFC 0009 (#13763, docs sweep and RFC status): keep the status header and the drift fixes. The RFC now says that hip-kernel-provider exports the serialization hooks. Testing (gfx942 MI300X, image hipdnn_20261017_gfx942): - Baseline, clean develop 0bb1c36: build OK, quick set 62/62 (job 68147020). - Merged tree: build OK, quick set 62/62 (job 68147243). - Targeted provider tests: unit 195/195, IntegrationGpuKernelIngestor* 26/26 (job 68147244). Plugin SDK targeted tests 309/309 (job 68147245). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Conflict in .pre-commit-config.yaml (#11659 expanded the comment on the reference-data exclusion): keep the generated ingestor_plan_generated.h exclusion and the expanded comment. No exclusion pattern changes. Committed on the user's request while gfx942 testing (full build, quick set, provider unit and integration tests; job 68147835) is still running. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
… code Every claim was checked against the code. Corrections: - Last verified names develop 0bb1c36 with #13315 applied. - Plugin payload contract: the bytes stay valid until the destroy hook, hipDNN refuses an empty payload, and destroy runs once after each successful serialize and never after a failed one. - The build calculates the schema SHA-256; a unit test compares it. - Restore check 11: the stored kpack target is the base architecture of the GPU that saved the plan; a generic archive entry does not widen it. - One save refusal (a code object too large for the format) uses the load prefix. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Contributor
|
blocked by unrelated test failures, issuing override merge. |
adickin-amd
disabled auto-merge
October 23, 2026 22:04
adickin-amd
deleted the
users/cmiserva/almiopen-2642/save-execution-plans
branch
October 23, 2026 22:04
BradPepersAMD
added a commit
that referenced
this pull request
Oct 24, 2026
…-saving check develop's #13315 checks that every ingestor pack's handler reports supportsSaving() exactly when it overrides save and restore, and asserts its symbol list covers every pack. With FlyDSL compiled in, the pack table has eight more packs. Their handlers do not save plans yet (the engines simply do not report SUPPORTS_EXECUTION_PLAN_SERIALIZATION), which the check now confirms for each. JIRA ID: ALMIOPEN-2953 Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
19 of 20 tasks
adickin-amd
pushed a commit
that referenced
this pull request
Oct 26, 2026
…tor tests (#14412) ## Summary Fixes two hip-kernel-provider tests that fail on develop since the gfx1151 split-KV decode engine (#13897) landed beside saving ingestor execution plans (#13315). Neither PR's CI ran with the other. Test-only change. JIRA ID : ALMIOPEN-2642 JIRA ID : AICK-1847 ## Risk Assessment Risk 1 of 5. Two test files change; no product code, descriptors or build wiring. ## ASIC Coverage No ASIC impact: both tests are host-only (a gtest that does no GPU work, and a pytest over descriptor JSON). Passing PR CI is sufficient. ## Testing Summary - `TestKernelIngestorEngine.EveryHandlerReportsSupportsSavingExactlyWhenItSavesAndRestores` fails on develop (4 dispatch symbols listed, 5 packs registered) and passes with the fix. The decode pack's handler is probed as one that cannot save. - `test_real_attention_dense_bundles_launch_with_the_ingestor_geometry` fails on develop with `KeyError: 'kernel_source'` and passes with the fix, with both parameters skipped, as before #13897. The test finds bundles one directory level deep, and the only `build_attention_dense` bundle (gfx950) keeps its KDP one level deeper in its DVC output, so the check has no kernels to compare on the shipped descriptors. Making it reach that bundle is a separate change. ## Testing Checklist - [x] Handler save-capability test - `hip_kernel_provider_tests --gtest_filter='TestKernelIngestorEngine.EveryHandlerReportsSupportsSavingExactlyWhenItSavesAndRestores'` (GPU hidden) - Status: Passed (1/1; fails on develop) - [x] rocKE corpus geometry test - `pytest descriptor-packaging/tests -m quick -k test_real_attention_dense_bundles_launch_with_the_ingestor_geometry` - Status: Passed, 2 skipped (fails on develop) - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - `TestKernelIngestorEngine.cpp`: `PACK_DISPATCH_SYMBOLS` gains `hipkernel.gfx1151_wmma_attention_decode.dispatch`, so it again lists one dispatch symbol per entry of `ingestorPacks()`. - `test_desk_check_rocke_corpus.py`: reads `kernel_source` with `.get()`, so a multi-launch kernel, which carries `sources` instead, drops out of the `build_attention_dense` kernels the test checks. Related: #13897, #13315.
4 of 5 tasks
This was referenced Oct 28, 2026
GarbyGilbertAMD
added a commit
that referenced
this pull request
Oct 30, 2026
… launch-set checks and an end-to-end test (#14215) ## Summary Follow-ups to #14158 (multi-launch `sources` kernel descriptors). A `sources` map may now author `hip` and `hsaco` entries as well as `rocke`, so a HIP or prebuilt-code-object multi-stage kernel can go through the packer. The packer refuses a pack in which kernels served by one dispatch handler name different launch sets. A device test runs a packer-built multi-launch kernel end to end through an engine, with a workspace hand-off between its launches. Saving an execution plan of a multi-launch kernel is refused for having several launches, rather than as an `embedded_source` kernel. The remaining review items from #14158 are closed with tests, tooling fixes and docs. JIRA ID : ALMIOPEN-2973 ## Risk Assessment Risk 3 of 5. The change extends the packer (`hkp_pack`) and IngestorGenerator with a new authoring path, adds a pack-time refusal, and adds three inline helpers to the plugin SDK ingestor headers. Single-launch descriptors and rocke `sources` pack and load as before; the state manager's launch-set and adaptability rules are unchanged and only move into shared helpers. The new packer refusal can fail a pack that used to succeed, but such a pack already had kernels dropped at load. Execution-plan saving refuses a multi-launch kernel with the same status as before and a corrected message, and the serialization behavior note is unchanged for every engine. Consumer-agreement evidence stays rocke-only, so hip/hsaco `sources` kernels are treated like their single-launch counterparts (a declaration with `metadata_fields` over them is refused). The packer, the state manager and the validator also overlap with the open kernel-slots implementation (#13972), so whichever lands second will need a rebase. ## ASIC Coverage Host-side packaging, loading, plan-save checks and tooling, plus a test-only fixture. No kernel, support surface, selection rule or default changes. Passing PR CI is sufficient. The device tests also ran and passed on gfx90a. ## Testing Summary - End to end on device: `TestMultiLaunchEndToEnd` loads a UKD 1.1 hip `sources` kernel packed by `hkp_pack` (two launches on one translation unit, one code object, two symbols), selects it through an engine with a test-only dispatch handler, and executes it. The first launch writes the workspace and the second reads it. The asserted output is reachable only if both launches run, in order, over one workspace. Host-only cases check that a handler declaring fewer launches drops the kernel and logs both launch sets. - Plan saving: a plan over a multi-launch kernel whose launches are all kpack is refused for having several launches, and the message does not call it an `embedded_source` kernel. - Packer: hip and hsaco `sources` (compile and read once per shared source, per-launch symbol and signature, provenance, desk check, refusals for `role`, a missing hsaco arch, mixed kinds and `metadata_fields`). Also the pack-time launch-set check, including two KDPs that share one UDD and two UDDs that name one dispatch symbol, and kills for previously surviving mutants: compile grouping by library and toc key, the desk-check reader per launch, and launch-order insensitivity. - Probes: hip and hsaco launches checked by archive membership, digest, symbol and signature without rocke evidence; rocke evidence still required. Also per-launch sweep keys and group keys, and launch-order insensitivity. - IngestorGenerator: the generated multi-launch `launchNames()` stub is now compiled, its census arm parsed, and its launch set checked against the config. `launches` on hip and hsaco kernel sources is accepted. `verify_variant_sets` keys hip and hsaco launches by their own identity and ignores `sources` key order. - Validator: a committed fixture whose kernels name different launch sets fails validation and names both sets. - Mixed-version check (manual): a validator built from develop just before #14158 skips UKD 1.1 kernels with a version warning instead of misloading them. A pack holding only 1.1 kernels is dropped; in a mixed pack its 1.0 kernel is kept. - Mutation checks: every new test was run against a targeted source mutant and failed. Examples: state manager keeping every kernel, plan passing no workspace, every launch given the first launch's symbol, signatures keyed per code object instead of per symbol, plan capture reading the kernel's `source` instead of its launches. - Review follow-ups from #14158. The new tests cover: - each launch's symbol, toc key and sha256 changing the engine-selector revision; - a device-free check that each launch's signature is checked against its own expected signature before its archive is opened; - the desk check catching a toc key shared between a single-launch kernel and one launch of a multi-launch kernel; - an accessor that raises AttributeError failing the observation instead of counting as unreported; - empty launch `target`/`role` overrides being refused; - a nesting case that differs only in a launch name. Each was run against a targeted source mutant and failed. ## Testing Checklist - [x] hipDNN superbuild, Linux, kernel ingestor, tooling tests and rocKE ON, GPU hidden - `ctest --test-dir build --output-on-failure --no-tests=error` - Status: Passed (236 registered; the 3 skips are 2 sample tests and the rocKE arch-domain reproduction test, none of them changed here) - [x] Multi-launch and plan-save device tests - `hip_kernel_provider_tests --gtest_filter='TestMultiLaunch*:TestLaunchKernelDefinition*:TestIngestorPlanCapture*:TestRestoredIngestorPlan*:TestIngestorPlanRoundTrip*'` - Devices: gfx90a - Status: Passed (26/26) - [x] IngestorGenerator suite - `ctest -R hipdnn_ingestor_generator_pytest` (strict skips) - Status: Passed (1545 passed, 0 skipped) - [x] Packer, probe and validator suites - `ctest -R 'hkp|probe|validate_descriptors'` - Status: Passed - [x] pre-commit - GitHub check - Status: Passed - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - **`hkp_pack`:** authored `sources` entries may be `hip` (`source`, `entry`, `build`) or `hsaco` (`file`, `symbol`), one kind per map. Entries sharing a translation unit and build, or an hsaco file, share one code object and toc key; each launch keeps its own symbol and signature. Provenance records each launch's authored identity under `provenance.sources`. `role`, consumer agreement and `check_claims` stay rocke-only; `embedded_source` and authored `kpack` entries stay refused. - **Pack-time launch-set check:** within an arch shard, every kernel served by one dispatch symbol (the `dispatch_symbol` of its KDP's UDD, which is how the loader resolves the handler) must name one launch set (a single-launch kernel's set is empty); otherwise the pack fails naming the symbol, each KDP and each kernel's set. KDPs that share a UDD, or name one symbol from two UDDs, are held to one set. The check sees one packaging root, so two roots that register the same symbol are not compared. The probe end-to-end fixture's fused multi-launch KDP gets a UDD of its own, since its single-launch siblings share the original one. - **Plugin SDK:** `Descriptors.hpp` holds the readers of a kernel's launches. `launchNamesOf()` builds a kernel's launch-name set, for the state manager and `hipdnn_validate_descriptors`. A multi-launch kernel leaves `source` default, which reads as an `embedded_source` kernel, so `singleLaunchSource()` returns a single-launch kernel's one source and null for a multi-launch kernel, and `allLaunchSources()` tests every launch. The state manager's adaptability check, `buildForSourceKind`, plan capture and the serialization behavior note read through them; the engine-selector digest still hashes both fields as before. - **Execution-plan saving:** the saved payload holds one kernel image, so capture refuses a multi-launch kernel for having several launches instead of reading its default `source`. Saving these plans is tracked in ALMIOPEN-3031. - **Probes:** `Kind.launch_evidence` decides whether a kind's launches carry rocke effective-spec evidence; hip and hsaco launches get the structural per-launch checks. - **IngestorGenerator:** `launches` on hip and hsaco kernel sources, with per-kind override keys; `verify_variant_sets` identities for hip and hsaco launches. - **Tests and fixtures:** new `multi_launch` descriptor set, packed beside the unit and integration roots so no Container loads it; validator fixture `multi_launch_mismatch`; the `_restamp` test helper moved to `pack_helpers.py`. - **Docs:** RFC 0017 §15.4/§15.5 corrected and the accepted `sources` kinds stated; RFC 0009, the provider README and the plugin development guide state that a multi-launch kernel's plan cannot be saved yet; a worked multi-launch handler in the ingestor-engine `native-pack.md` guide; packaging, generator and fixture READMEs updated. - **Review follow-ups (#14158):** - The desk check's toc-key uniqueness now counts two kernels as one when any launch of one shares a toc key with any launch of the other. - Unknown `sources` kinds share one refusal message. - Empty launch `target`/`role` overrides are refused. - `launchKernelDefinition` no longer copies the `sources` map only to discard it. - The RFC 0017 §15.1 example is marked abridged, and the "fill every Launch" loader check is described as arriving with `launches[]`. Related: #14158 (multi-launch kernel descriptors), #13315 (saving and loading ingestor execution plans), #13839 (RFC 0021, kernel slots), #13972 (kernel-slots implementation), #13897 (RDNA SDPA catalog, a future `sources` user).
GarbyGilbertAMD
added a commit
that referenced
this pull request
Oct 30, 2026
…oaded from (#14517) ## Summary A plan loaded from saved bytes could not be saved again: the backend asked the plugin to serialize the restored execution context, and hip-kernel-provider refuses to capture a restored plan, so `to_binary()` of a graph loaded with its plan and `to_compiled_plan_binary()` of a loaded plan both failed. The backend now keeps the plan bytes a plan was loaded from and returns them on a later save, without calling the plugin. When the input is a graph-and-plan container, only its plan part is kept. The design decision is recorded in RFC 0009 ("Decision: Re-Save From the Loaded Bytes"). JIRA ID : ALMIOPEN-2940 Follow-up to #13315, which deferred re-saving a loaded plan to this ticket. ## Risk Assessment 2 — `ExecutionPlanDescriptor::deserializeBackendPlan` seeds the existing serialized-plan cache with one added line, and passes the plugin payload to the plugin without a temporary copy; the save path for a built plan is unchanged. No schema, public API or plugin ABI change. Behaviour change: `to_binary()` of a graph loaded with its plan and `to_compiled_plan_binary()` of a loaded plan now succeed for every plugin that exports the three serialization hooks, and hipDNN no longer calls the plugin's save hook for a loaded plan. A re-save returns the plan bytes the plan was loaded from, so it keeps the original provider and format versions; a newer kernel in an upgraded provider never enters it (rebuild the plan from its graph to pick one up). Memory cost: every loaded plan keeps one copy of its plan bytes for its lifetime, whether or not it is saved again, while a built plan keeps a copy only after its first save. For hip-kernel-provider plans the bytes include the kernel code object. Load no longer copies the plugin payload before it hands it to the plugin, which lowers peak memory at load. ## ASIC Coverage Architecture-independent host code in the backend descriptor; the provider change is a message and comments. PR CI is sufficient. Verified locally on gfx1152 (Windows), where all four GPU save/load cases ran; PR CI provides the Linux build and the other architectures. ## Testing Summary - Backend unit: a loaded plan re-saves byte-identically without calling the plugin; a version 1 envelope stays at version 1 (a rebuilt envelope would claim version 2 with no tensor alignments and fail to load); a failed restore leaves the plan unfinalized and unserializable; a built plan still fills the cache on its first save, including concurrent saves and a retry after a failed fill. - Provider unit: the save hook still refuses a direct call on a restored plan, with the reworded message. - GPU integration: pointwise and convolution plans, each as a graph with its plan and as a plan alone, are saved, loaded in a fresh handle, saved again, loaded a second time in a fresh handle and validated. For a graph with its plan, the plan bytes that the loaded and the reloaded graph save equal the plan bytes of the original plan. - Mutation check: with the cache fill removed, the new backend tests and all four GPU save/load cases fail. ## Testing Checklist - [x] hipDNN backend unit tests - `hipdnn_backend_tests` (`-DHIPDNN_ENABLE_KERNEL_INGESTOR=ON`) - Status: Passed - [x] hip-kernel-provider unit tests - `hip_kernel_provider_tests` - Status: Passed - [x] Save/load GPU integration - `hip_kernel_provider_integration_tests --gtest_filter=IntegrationGpuKernelIngestorSaveLoad*` - ASIC: gfx1152 - Status: Passed - [x] Test-name validation - `ctest -R test_name_validation` (hipDNN and hip-kernel-provider) - Status: Passed - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - `ExecutionPlanDescriptor::deserializeBackendPlan` assigns the plan bytes it was loaded from to `_serializedPlanCache` after the plugin restores the context and before `finalize()`; `serializeBackendPlan` returns the cache when it is populated. - hip-kernel-provider keeps refusing a direct save-hook call on a restored plan; the message and comments state that hipDNN saves a loaded plan from its plan bytes and never makes that call. - Backend descriptor tests and the GPU save/load integration tests cover the re-save. - `Graph::serialize_compiled_plan` documents that a plan loaded from bytes can be saved. - RFC 0009, `PluginDevelopment.md` and the hip-kernel-provider README describe the behaviour and the memory cost.
13 of 21 tasks
AviralGoelAMD
added a commit
that referenced
this pull request
Nov 2, 2026
…r the develop merge rocKE #14367 changed _num_segments on gfx942: a head-size-256, block-16 decode rule (about 2 CTAs per CU, at most max_seqlen_k // 32 segments, at least 2), and the D128 pre-bump clamp now counts the decode grid (num_seqs q-blocks). gfx942SplitRuleSegments restates both; the split formula is factored out so the pre-bump can take the decode-grid num_2d. The split-rule table is regenerated and its grid widened (review): GQA ratios 2 and 8, and pinned rows for the Skv 8192 decode-rule threshold, the D128 pre-bump Skv 32768 bound, the D256 block-16 floor and cap, and a qpk-8 row the grid cannot otherwise distinguish. 8,650 -> 14,419 rows; new branch anchors. Develop API and wiring updates the merge left: - TestGfx942AttentionUnified3dDispatch: unsortedDefinitions() returns pointers (#13207); copy the kernel out before the state manager goes. - The gfx942 integration lane reserves its TOML with the census shard, as the gfx950 lane does (#14111 check). - TestKernelIngestorEngine lists the gfx942 pack's dispatch symbol (#13315).
cgmillette
added a commit
that referenced
this pull request
Nov 4, 2026
…r engines (#14242) ## Summary Packages rocKE's gfx950 unified attention as two hipDNN kernel-ingestor engines: `hipkernel:Gfx950AttentionUnified2D` (prefill and the decode shapes 3D does not cover) and `hipkernel:Gfx950AttentionUnified3D` (split-KV decode, a segment and a reduce launch per kernel through #14158's multi-launch `sources`). Both ship AOT, so an SDPA-forward graph on gfx950 can be served by unified attention with no JIT. Both catalogs are reduced to what a correctness-checked MI355X sweep measured as needed: Unified2D ships 222 of its 792-kernel full space and Unified3D 123 of 348, with no graph losing coverage or more than the measured noise (the final reduced-vs-full sweep is below). JIRA ID : AICK-2542 JIRA ID : AICK-2682 ## Risk Assessment Risk 2 of 5. The PR adds two engines and their descriptors. No existing engine's code, catalog or ranking changes: an unpinned SdpaFwd graph still ranks `hipkernel:Gfx950AttentionDense` first (device test `DenseStillRanksAheadWhenUnpinned`). The shared pieces it touches: - generator tooling: `registry_enumerator` full-space and multi-launch sections, the regeneration gates, provenance stamping, and two fixes in `variant_reduce`. Only `provenance.generated_from` is opt-in. The expanded registry output now always writes per-kernel `provenance.rocke_spec_id` / `rocke_tuning_id`, including Dense's 840 catalog kernels (its manifest pins a `spec_id`). The compact `variants:` selections drop it, so no shipped descriptor changes, but any other consumer of the expanded output sees the new field; - one rocKE builder, `build_unified_attention_3d_tiled_module`; - the ingestor save census's list of dispatch symbols. These are covered by their own tests and by the existing dense regeneration gates, which still pass. 3D is the first production user of `launchNames()` (#14158). ## ASIC Coverage gfx950 only: both engines ship only for gfx950, and other architectures' shards carry neither engine. PR CI has no gfx950 device lane (`hkp-probe-gfx950` packs and asserts but launches nothing), so the device evidence comes from MI355X (and, before the rebase, MI350X) runs. ## Testing Summary - Host: unit and census tests for both packs, the arch-neutral geometry, and the generator suite under `hipdnn_strict_skips`, including regeneration gates g1–g7 for all three gfx950 manifests; static descriptor gates on the installed shards (variant sets, coverage rungs 1–2, validator, launch surface). - Device, MI355X (gfx950), on the shipped (reduced) catalogs at this head's code: - engine-pinned SdpaFwd quick and standard suites against the CPU and GPU references, both engines plus a Dense regression run; - the knobs suite: every shipped kernel class forced and checked against the CPU reference, plus declines, precedence and a NaN-prefilled workspace; - a timing comparison of every engine that serves each corpus graph, and PyTorch; - RUNBOOK stage 7's final sweep: the full catalog against the reduced one in one session. - How the catalogs were reduced: a 3-round dnn-benchmark sweep with correctness on, over 188 prefill graphs (2D) and 88 decode graphs (3D), then `variant_reduce`. ## Packaging and performance statistics All numbers below come from tool output (packaging logs, device suite runs, dnn-benchmark sweeps). **Packaging** (gfx950 shard, shipped catalogs; pack time is the build's own `hkp_pack` command per engine, median of 3, relative to Dense's on the same job) | Engine | Kernels (full space) | AOT budget | Code objects, zstd | KDP | Pack time vs Dense, build default workers | Serial time per kernel vs Dense | |---|---:|---:|---:|---:|---:|---:| | Unified2D | 222 (792) | 1000 | 1.55 MiB | 0.29 MB | 0.11× | 1.58× | | Unified3D | 123 (348), 246 launches | 400 | 0.64 MiB | 0.26 MB | 0.41× | 1.06× | | Dense (catalog unchanged, for scale) | 840 (52 distinct binaries since #14059) | 1000 | 6.16 MiB | 1.26 MB | 1 | 1 | - The gfx950 kpack is 8.49 MiB at this head; the two engines are 2.19 MiB, 26% of it. Since #14059 (runtime head counts for Dense), Dense's 840 catalog kernels compile to 52 distinct binaries, but the kpack still stores one frame per kernel, so Dense grew from 5.79 to 6.16 MiB. Collapsing those entries is #14059's stated follow-up. The pack-time ratios and the full-space figure (14.45 MiB) were measured against Dense before #14059. - Provider load (`hipdnn_list_engines`, median of 5) is 1.24× its time without these descriptors. **Correctness and coverage** (MI355X, shipped catalogs) | Suite | Unified2D | Unified3D | Dense | |---|---|---|---| | SdpaFwd quick + standard, CPU and GPU reference (served / declined / failed) | 43 / 6 / 0 | 4 / 45 / 0 | 45 / 4 / 0 | | Knobs suite (forced kernel classes, declines, precedence) | 30 / 30 (both engines) | | 62 / 62 | | Corpus graphs served and validated against PyTorch | 188 / 188 | 82 / 88 | | Every decline is intentional: Unified2D leaves decode to Unified3D and declines Hkv > Hq; Unified3D serves decode only, and its 6 corpus declines are batch-256 decode graphs whose K/V is over the 32-bit addressing cap (ASM SDPA serves those). **Kernel time** (MI355X, autotuned, median of 3 rounds; ratio = Unified / other, below 1 = Unified faster; every engine's round-to-round drift ≤ 0.4%). The Unified2D rows were re-measured at this head, on the kernels develop's tiled-2D changes recompiled (see the checklist). | | Graphs both serve | Geomean ratio | Time-weighted (other / Unified) | Unified faster on | |---|---:|---:|---:|---:| | Unified3D vs ASM SDPA (decode) | 60 | 0.200 | 8.17 | 60 / 60 | | Unified3D vs PyTorch | 82 | 0.062 | 34.2 | 82 / 82 | | Unified2D vs Dense (prefill) | 179 | 2.37 | 0.38 | 0 / 179 | | Unified2D vs PyTorch | 188 | 1.38 | 0.55 | 39 / 188 | - Unified3D serves 22 decode graphs no existing hipDNN engine serves. - Unified2D is slower than Dense wherever Dense serves, and Dense keeps ranking first there; it adds coverage Dense lacks (the 9 bottom-right causal graphs in these corpora). **Variant reduction** | Engine | Kernels | Chosen eps | Measured drift | Final sweep, reduced / full kernel time (geomean, rounds 1–3) | |---|---|---:|---:|---| | Unified2D | 792 → 222 | 0.02 | 0.62% | 1.000, 0.999, 1.000 (u2d_attn); 1.000, 0.999, 1.000 (u2d_rocke) | | Unified3D | 348 → 123 | 0.03 | 1.08% | 1.004, 1.002, 0.999 (u3d_attn); 1.002, 1.007, 1.004 (u3d_rocke) | Both arms of the final sweep served the same graphs and passed correctness on every corpus. ## Testing Checklist - [x] Unit (host, no GPU) - `HIP_VISIBLE_DEVICES=-1 hip_kernel_provider_tests` - 1454 passed, including `TestKernelIngestorEngine.EveryHandlerReportsSupportsSavingExactlyWhenItSavesAndRestores`. Our build host also fails 11 Dense signature/dispatch tests: it has a `/nonexistent` directory, which those tests assume is absent. They pass in CI - Status: Passed - [x] Census - `ctest -R 'hkp-census|rocke_unified'` - 36/36 (the census inventories are the reduced catalogs') - Status: Passed - [x] Generator suite (strict skips) - `pytest -p hipdnn_strict_skips IngestorGenerator/tests` - 1646 passed, 0 skipped - Status: Passed - [x] Regeneration gates on the reduced selections - `tests/test_registry_regeneration.py` (g1–g7) 39 passed; AOT budget tests 12 passed - Status: Passed - [x] rocKE attention CPU tests and packaging guards - 2D/3D mask_type, 3D module, decode split, KV cache policy, spec-predicate parity: 145 passed, plus 6 develop KV-cache tests that need the compiled `rocke_engine`; `test_hkp_pack_producer_guards.py` and `test_gfx950_attention_unified_3d_bundle_coverage.py`: 26 passed - Status: Passed - [x] Golden IR, before the reduction - `ctest -R rocke_golden_static` - the reduction changes no rocKE code - Status: Passed - [x] Launch surface - `launch_surface.py configs/gfx950_attention_unified_{2d,3d}.engine.yaml --check` - Status: Passed - [x] Variant sets - `verify_variant_sets.py --mode structural` and `--mode full --arch gfx950 --manifest …` - 2D 222/222 distinct binaries, 3D 123/123 - Status: Passed - [x] Coverage gate rungs 1–2 - `coverage_gate.py --mode full --expect-engine …` - Status: Passed - [x] Validator - `hipdnn_validate_descriptors <shard> --expect-engine … --json` - success, no warnings - Status: Passed - [x] Packaging probe - gfx942 configure with `HIPKERNELPROVIDER_ENABLE_PACKAGING_PROBES=ON`, `ctest -R '^hkp-probe-gfx950'` - 12/12; `hkp-probe-gfx950` packs 51 UKDs of the current descriptors in a gfx942-only build - Status: Passed - [x] Device suites, MI355X - engine-pinned SdpaFwd quick + standard with GPU reference (`hip_kernel_provider_gfx950_attention_*_gpu_ref_integration_tests`) and CPU reference (`*-external-integration_{quick,standard}_suite`) - 0 failures: Unified2D 43 served / 6 declined, Unified3D 4 / 45, Dense 45 / 4 - Status: Passed - [x] Device knobs suite, MI355X - `hip_kernel_provider_integration_tests --gtest_filter='*Gfx950AttentionUnified*'` - 30/30; Dense knobs 62/62 - Status: Passed - [x] Fault re-check, MI355X - the whole `u2d_rocke` corpus in one autotuned process, twice and then with serialized launches - 93/93 each, 0 errors - Status: Passed - [x] Reduction sweep + `variant_reduce.py` - 3 rounds, correctness on; 2D `SWEEP_DONE`, 3D `SWEEP_DONE`; 2D 792 → 222, 3D 348 → 123 - Status: Passed - [x] Final sweep, reduced against full (RUNBOOK §7) - same coverage, correctness passed on both arms, ratios above - Status: Passed - [x] Unified3D catalog regenerated for #14367's LLC-aware K/V cache policy - all 123 kernels now `kv_cache_policy: "all"` (was `"stream"`), same kernel classes; MI355X re-run: suites and knobs unchanged, 246/246 decode rows correct, kernel time against the `"stream"` catalog geomean 0.998, time-weighted 0.999 - Status: Passed - [x] Device evidence at this head (rebased onto develop 2026-11-03, after the review-round refactors) - Dense's and Unified3D's code objects are byte-identical to the previous device runs; develop's #13652 recompiled 201 of the 222 Unified2D code objects (same kernels and ids), and the Unified2D/3D matchers moved onto a shared header, so both engines' device runs were repeated at this head: suites Unified2D 43 / 6 / 0, Unified3D 4 / 45 / 0, Dense 45 / 4 / 0 (passed / skipped / failed), knobs 30/30 and 62/62; compare 188/188 Unified2D graphs (564/564 rows) and 82 Unified3D graphs (246/246 rows) correct against PyTorch in all 3 rounds. Kernel time vs the previous runs: Unified2D geomean 0.985 (none of 188 graphs >2% slower), Unified3D 1.003 (ASM_SDPA control 1.006). A final rebase picked up #14059, which recompiled Dense's 840 code objects only (Unified2D/3D byte-identical to the run above); the suites and knobs were re-run at this head with the same results, `DenseStillRanksAheadWhenUnpinned` included - Status: Passed - [x] DVC - both reduced KDPs on the `ingestor` remote; `dvc status -c` in sync - Status: Passed - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - **Unified2D pack** (`Gfx950AttentionUnified2DNative.cpp`): matches BSHD SDPA forward on the identity paged-KV layout and builds block_tables, seq_lens and query_start_len in the workspace. - It declines: batch > 511 (the baked 9-iteration binary search), Skv not a multiple of 16, i32 KV over 0x7FFF0000, any window, and the decode cells Unified3D covers (Sq = 1 with the GQA group dividing 16). - Engine knobs: tile_size, num_warps, block_m_per_warp, use_mfma_32x32 and use_transposed_qk_32x32. - **Unified3D pack** (`Gfx950AttentionUnified3DNative.cpp`): decode only. - `launchNames()` is `{segment, reduce}`, and both definitions are resolved through `launchKernelDefinition` with a signature check before any load. - The segment launch writes `[token, Hq, seg, D]` scratch and the reduce always runs, so the workspace needs no zeroing. - Knob: num_segments. No batch limit. - **Arch-neutral geometry** (`UnifiedAttentionGeometry.hpp`): the launch geometry and workspace layouts take a `UnifiedAttentionArchTraits`, and the gfx950 headers are thin wrappers, so a gfx942 pack (AICK-2544) only supplies traits. - **rocKE**: `build_unified_attention_3d_tiled_module` packs the segment and its derived reduce (`reduce_spec_for`) as one `KernelModule`. The module predicate gates both halves. The 3D pack compiles develop's running-max fix for split-KV (#14276). - **Generator**: - `enumerate.also` sections may use `space: full` and carry manifest `launches`. The enumerator refuses a registry that pairs a reduce when the manifest declares no launches, and keeps one default per cell. - The gates are launch-aware: g1 keys on launches, g3 uses `--subset` for full spaces, g4 compares launch sets, and g6 walks every launch and checks module roles. - g2 and g3 run per engine with the shipped ids reused (#14109), comparing every descriptor byte for byte. - `generate.py --stamp-provenance` stamps `provenance.generated_from` (config, manifest, generator and rocKE commits). It is opt-in, so a default regeneration stays byte-identical. - `variant_reduce` fixes: - it reads the format-2 winner lines the runtime has written since #13028; - it keys a graph on its content with `Graph.id` (minted fresh at each finalization) zeroed, so sweep rounds match. - **Catalogs** (the manifests enumerate the full spaces; the selections are `variant_reduce`'s): - Unified2D: 222 kernels over 198 cells (33 head configs × bf16/fp16 × 3 masks): the narrow nw4/mw16/t64 geometry in every cell, transposed 32x32 nw4/mw32/t64 in 21 and narrow nw2/mw16/t32 in 3. The wide 32x32 geometry won no graph and is not shipped. - Unified3D: 123 kernels: num_segments 16 in all 116 decode cells, 32 or 64 in 7 bf16 bottom-right D128 variants. The K/V cache policy follows #14367's LLC-aware rule, which picks `"all"` for every shipped decode cell. - **Descriptors**: both are generated with develop's generator (post-#13028 UED/UHD, UKD `sources` for 3D), every kernel id reused from the full catalog. A regeneration with the shipped ids reused reproduces them byte for byte. The KDPs live in DVC (`ingestor` remote). - **Shared code** (one home per fact): - `packs/UnifiedAttentionGraphMatch.hpp`: the graph-match helpers both unified packs use (operand/BSHD checks, SDPA node, mask classification, dtype, metadata readers); the tests' graph builder is `UnifiedAttentionTestGraphs.hpp`. - rocKE `kernels/common/attention_mask_type.py` (mask ordinals and band -> kind, used by both tiled specs and re-exported by dispatch) and `attention_split_kv_reduce.py` (the segment -> reduce derivation the dispatcher and the module export share). - CMake registers Dense, Unified2D and Unified3D by one loop per lane (external integration, gpu_ref, census). - **Tests**: - The device knobs suite (`IntegrationGpuGfx950AttentionUnifiedKnobs.cpp`) adds a harness `prepareWorkspace` hook. It forces the default geometry/segment count across the head configurations and each other shipped class on a cell that carries it, and `CoversEveryCatalogKernelClass` holds those tables to the catalog. - The engine-pinned `gpu_ref` targets and external-integration lanes cover both engines. - The ingestor save census names all seven pack dispatch symbols. - The Unified2D/3D integration TOMLs are reserved with the gfx950 census shard (`hkp_reserve_census_shard(... FILES ...)`, as #14573 requires of every shard entry). **Constraints and follow-ups** - #13372 and #13889 are both in develop; this branch no longer carries a copy of #13889. - Heads are catalog cells (baked), window = 0, 2D batch ≤ 511, i32 KV ≤ 0x7FFF0000. D256 is deferred for both engines. The kernarg work is tracked in AICK-2675. - There is no 2D/3D `select_path` crossover replay: the split is the static decode rule above. - Runtimes older than UKD 1.1 skip the 3D descriptors. Strided-KV 3D is a follow-up. - Follow-up: execution-plan save/restore (#13315) for both handlers. Today they report `supportsSaving() == false`, as the RDNA handler does. - One earlier MI355X sweep hit a single GPU fault in Unified2D on `u2d_rocke/rocke_sdpa_035` (bf16, D 128, Hq 64 / Hkv 8, causal, S 512). It has not reproduced since: not in the forced geometries on that cell, 15 single-graph runs, 4 full-corpus runs on MI350X, or the MI355X full-corpus re-check above. - Unified2D kernel time moved with develop's tiled-2D changes (#14542, #14616), not with this PR: at this head it is 1.035× (geomean) its pre-rebase time on the same 188 graphs, up to 1.36× on long causal bf16 D128 prefill (S 8192), while Dense in the same runs is flat (0.997). Dense still ranks first wherever it serves; the table above is at this head. - Measurement note: the pinned dnn-benchmarking (v1 result schema, which `sweep.py` reads) has no bottom-right causal PyTorch reference. The sweeps used upstream dnn-benchmarking #84's reference, with the 9 bottom-right corpus graphs written as `causal_mask: false, right_bound: 0, diagonal_alignment: BOTTOM_RIGHT` (the same band under hipDNN's `resolveDiagonalBand`): #84 reads `causal_mask: true` as top-left regardless of the alignment. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Christopher Millette <63608002+cgmillette@users.noreply.github.com>
cycheng
pushed a commit
that referenced
this pull request
Nov 6, 2026
…r engines (#14242) ## Summary Packages rocKE's gfx950 unified attention as two hipDNN kernel-ingestor engines: `hipkernel:Gfx950AttentionUnified2D` (prefill and the decode shapes 3D does not cover) and `hipkernel:Gfx950AttentionUnified3D` (split-KV decode, a segment and a reduce launch per kernel through #14158's multi-launch `sources`). Both ship AOT, so an SDPA-forward graph on gfx950 can be served by unified attention with no JIT. Both catalogs are reduced to what a correctness-checked MI355X sweep measured as needed: Unified2D ships 222 of its 792-kernel full space and Unified3D 123 of 348, with no graph losing coverage or more than the measured noise (the final reduced-vs-full sweep is below). JIRA ID : AICK-2542 JIRA ID : AICK-2682 ## Risk Assessment Risk 2 of 5. The PR adds two engines and their descriptors. No existing engine's code, catalog or ranking changes: an unpinned SdpaFwd graph still ranks `hipkernel:Gfx950AttentionDense` first (device test `DenseStillRanksAheadWhenUnpinned`). The shared pieces it touches: - generator tooling: `registry_enumerator` full-space and multi-launch sections, the regeneration gates, provenance stamping, and two fixes in `variant_reduce`. Only `provenance.generated_from` is opt-in. The expanded registry output now always writes per-kernel `provenance.rocke_spec_id` / `rocke_tuning_id`, including Dense's 840 catalog kernels (its manifest pins a `spec_id`). The compact `variants:` selections drop it, so no shipped descriptor changes, but any other consumer of the expanded output sees the new field; - one rocKE builder, `build_unified_attention_3d_tiled_module`; - the ingestor save census's list of dispatch symbols. These are covered by their own tests and by the existing dense regeneration gates, which still pass. 3D is the first production user of `launchNames()` (#14158). ## ASIC Coverage gfx950 only: both engines ship only for gfx950, and other architectures' shards carry neither engine. PR CI has no gfx950 device lane (`hkp-probe-gfx950` packs and asserts but launches nothing), so the device evidence comes from MI355X (and, before the rebase, MI350X) runs. ## Testing Summary - Host: unit and census tests for both packs, the arch-neutral geometry, and the generator suite under `hipdnn_strict_skips`, including regeneration gates g1–g7 for all three gfx950 manifests; static descriptor gates on the installed shards (variant sets, coverage rungs 1–2, validator, launch surface). - Device, MI355X (gfx950), on the shipped (reduced) catalogs at this head's code: - engine-pinned SdpaFwd quick and standard suites against the CPU and GPU references, both engines plus a Dense regression run; - the knobs suite: every shipped kernel class forced and checked against the CPU reference, plus declines, precedence and a NaN-prefilled workspace; - a timing comparison of every engine that serves each corpus graph, and PyTorch; - RUNBOOK stage 7's final sweep: the full catalog against the reduced one in one session. - How the catalogs were reduced: a 3-round dnn-benchmark sweep with correctness on, over 188 prefill graphs (2D) and 88 decode graphs (3D), then `variant_reduce`. ## Packaging and performance statistics All numbers below come from tool output (packaging logs, device suite runs, dnn-benchmark sweeps). **Packaging** (gfx950 shard, shipped catalogs; pack time is the build's own `hkp_pack` command per engine, median of 3, relative to Dense's on the same job) | Engine | Kernels (full space) | AOT budget | Code objects, zstd | KDP | Pack time vs Dense, build default workers | Serial time per kernel vs Dense | |---|---:|---:|---:|---:|---:|---:| | Unified2D | 222 (792) | 1000 | 1.55 MiB | 0.29 MB | 0.11× | 1.58× | | Unified3D | 123 (348), 246 launches | 400 | 0.64 MiB | 0.26 MB | 0.41× | 1.06× | | Dense (catalog unchanged, for scale) | 840 (52 distinct binaries since #14059) | 1000 | 6.16 MiB | 1.26 MB | 1 | 1 | - The gfx950 kpack is 8.49 MiB at this head; the two engines are 2.19 MiB, 26% of it. Since #14059 (runtime head counts for Dense), Dense's 840 catalog kernels compile to 52 distinct binaries, but the kpack still stores one frame per kernel, so Dense grew from 5.79 to 6.16 MiB. Collapsing those entries is #14059's stated follow-up. The pack-time ratios and the full-space figure (14.45 MiB) were measured against Dense before #14059. - Provider load (`hipdnn_list_engines`, median of 5) is 1.24× its time without these descriptors. **Correctness and coverage** (MI355X, shipped catalogs) | Suite | Unified2D | Unified3D | Dense | |---|---|---|---| | SdpaFwd quick + standard, CPU and GPU reference (served / declined / failed) | 43 / 6 / 0 | 4 / 45 / 0 | 45 / 4 / 0 | | Knobs suite (forced kernel classes, declines, precedence) | 30 / 30 (both engines) | | 62 / 62 | | Corpus graphs served and validated against PyTorch | 188 / 188 | 82 / 88 | | Every decline is intentional: Unified2D leaves decode to Unified3D and declines Hkv > Hq; Unified3D serves decode only, and its 6 corpus declines are batch-256 decode graphs whose K/V is over the 32-bit addressing cap (ASM SDPA serves those). **Kernel time** (MI355X, autotuned, median of 3 rounds; ratio = Unified / other, below 1 = Unified faster; every engine's round-to-round drift ≤ 0.4%). The Unified2D rows were re-measured at this head, on the kernels develop's tiled-2D changes recompiled (see the checklist). | | Graphs both serve | Geomean ratio | Time-weighted (other / Unified) | Unified faster on | |---|---:|---:|---:|---:| | Unified3D vs ASM SDPA (decode) | 60 | 0.200 | 8.17 | 60 / 60 | | Unified3D vs PyTorch | 82 | 0.062 | 34.2 | 82 / 82 | | Unified2D vs Dense (prefill) | 179 | 2.37 | 0.38 | 0 / 179 | | Unified2D vs PyTorch | 188 | 1.38 | 0.55 | 39 / 188 | - Unified3D serves 22 decode graphs no existing hipDNN engine serves. - Unified2D is slower than Dense wherever Dense serves, and Dense keeps ranking first there; it adds coverage Dense lacks (the 9 bottom-right causal graphs in these corpora). **Variant reduction** | Engine | Kernels | Chosen eps | Measured drift | Final sweep, reduced / full kernel time (geomean, rounds 1–3) | |---|---|---:|---:|---| | Unified2D | 792 → 222 | 0.02 | 0.62% | 1.000, 0.999, 1.000 (u2d_attn); 1.000, 0.999, 1.000 (u2d_rocke) | | Unified3D | 348 → 123 | 0.03 | 1.08% | 1.004, 1.002, 0.999 (u3d_attn); 1.002, 1.007, 1.004 (u3d_rocke) | Both arms of the final sweep served the same graphs and passed correctness on every corpus. ## Testing Checklist - [x] Unit (host, no GPU) - `HIP_VISIBLE_DEVICES=-1 hip_kernel_provider_tests` - 1454 passed, including `TestKernelIngestorEngine.EveryHandlerReportsSupportsSavingExactlyWhenItSavesAndRestores`. Our build host also fails 11 Dense signature/dispatch tests: it has a `/nonexistent` directory, which those tests assume is absent. They pass in CI - Status: Passed - [x] Census - `ctest -R 'hkp-census|rocke_unified'` - 36/36 (the census inventories are the reduced catalogs') - Status: Passed - [x] Generator suite (strict skips) - `pytest -p hipdnn_strict_skips IngestorGenerator/tests` - 1646 passed, 0 skipped - Status: Passed - [x] Regeneration gates on the reduced selections - `tests/test_registry_regeneration.py` (g1–g7) 39 passed; AOT budget tests 12 passed - Status: Passed - [x] rocKE attention CPU tests and packaging guards - 2D/3D mask_type, 3D module, decode split, KV cache policy, spec-predicate parity: 145 passed, plus 6 develop KV-cache tests that need the compiled `rocke_engine`; `test_hkp_pack_producer_guards.py` and `test_gfx950_attention_unified_3d_bundle_coverage.py`: 26 passed - Status: Passed - [x] Golden IR, before the reduction - `ctest -R rocke_golden_static` - the reduction changes no rocKE code - Status: Passed - [x] Launch surface - `launch_surface.py configs/gfx950_attention_unified_{2d,3d}.engine.yaml --check` - Status: Passed - [x] Variant sets - `verify_variant_sets.py --mode structural` and `--mode full --arch gfx950 --manifest …` - 2D 222/222 distinct binaries, 3D 123/123 - Status: Passed - [x] Coverage gate rungs 1–2 - `coverage_gate.py --mode full --expect-engine …` - Status: Passed - [x] Validator - `hipdnn_validate_descriptors <shard> --expect-engine … --json` - success, no warnings - Status: Passed - [x] Packaging probe - gfx942 configure with `HIPKERNELPROVIDER_ENABLE_PACKAGING_PROBES=ON`, `ctest -R '^hkp-probe-gfx950'` - 12/12; `hkp-probe-gfx950` packs 51 UKDs of the current descriptors in a gfx942-only build - Status: Passed - [x] Device suites, MI355X - engine-pinned SdpaFwd quick + standard with GPU reference (`hip_kernel_provider_gfx950_attention_*_gpu_ref_integration_tests`) and CPU reference (`*-external-integration_{quick,standard}_suite`) - 0 failures: Unified2D 43 served / 6 declined, Unified3D 4 / 45, Dense 45 / 4 - Status: Passed - [x] Device knobs suite, MI355X - `hip_kernel_provider_integration_tests --gtest_filter='*Gfx950AttentionUnified*'` - 30/30; Dense knobs 62/62 - Status: Passed - [x] Fault re-check, MI355X - the whole `u2d_rocke` corpus in one autotuned process, twice and then with serialized launches - 93/93 each, 0 errors - Status: Passed - [x] Reduction sweep + `variant_reduce.py` - 3 rounds, correctness on; 2D `SWEEP_DONE`, 3D `SWEEP_DONE`; 2D 792 → 222, 3D 348 → 123 - Status: Passed - [x] Final sweep, reduced against full (RUNBOOK §7) - same coverage, correctness passed on both arms, ratios above - Status: Passed - [x] Unified3D catalog regenerated for #14367's LLC-aware K/V cache policy - all 123 kernels now `kv_cache_policy: "all"` (was `"stream"`), same kernel classes; MI355X re-run: suites and knobs unchanged, 246/246 decode rows correct, kernel time against the `"stream"` catalog geomean 0.998, time-weighted 0.999 - Status: Passed - [x] Device evidence at this head (rebased onto develop 2026-11-03, after the review-round refactors) - Dense's and Unified3D's code objects are byte-identical to the previous device runs; develop's #13652 recompiled 201 of the 222 Unified2D code objects (same kernels and ids), and the Unified2D/3D matchers moved onto a shared header, so both engines' device runs were repeated at this head: suites Unified2D 43 / 6 / 0, Unified3D 4 / 45 / 0, Dense 45 / 4 / 0 (passed / skipped / failed), knobs 30/30 and 62/62; compare 188/188 Unified2D graphs (564/564 rows) and 82 Unified3D graphs (246/246 rows) correct against PyTorch in all 3 rounds. Kernel time vs the previous runs: Unified2D geomean 0.985 (none of 188 graphs >2% slower), Unified3D 1.003 (ASM_SDPA control 1.006). A final rebase picked up #14059, which recompiled Dense's 840 code objects only (Unified2D/3D byte-identical to the run above); the suites and knobs were re-run at this head with the same results, `DenseStillRanksAheadWhenUnpinned` included - Status: Passed - [x] DVC - both reduced KDPs on the `ingestor` remote; `dvc status -c` in sync - Status: Passed - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - **Unified2D pack** (`Gfx950AttentionUnified2DNative.cpp`): matches BSHD SDPA forward on the identity paged-KV layout and builds block_tables, seq_lens and query_start_len in the workspace. - It declines: batch > 511 (the baked 9-iteration binary search), Skv not a multiple of 16, i32 KV over 0x7FFF0000, any window, and the decode cells Unified3D covers (Sq = 1 with the GQA group dividing 16). - Engine knobs: tile_size, num_warps, block_m_per_warp, use_mfma_32x32 and use_transposed_qk_32x32. - **Unified3D pack** (`Gfx950AttentionUnified3DNative.cpp`): decode only. - `launchNames()` is `{segment, reduce}`, and both definitions are resolved through `launchKernelDefinition` with a signature check before any load. - The segment launch writes `[token, Hq, seg, D]` scratch and the reduce always runs, so the workspace needs no zeroing. - Knob: num_segments. No batch limit. - **Arch-neutral geometry** (`UnifiedAttentionGeometry.hpp`): the launch geometry and workspace layouts take a `UnifiedAttentionArchTraits`, and the gfx950 headers are thin wrappers, so a gfx942 pack (AICK-2544) only supplies traits. - **rocKE**: `build_unified_attention_3d_tiled_module` packs the segment and its derived reduce (`reduce_spec_for`) as one `KernelModule`. The module predicate gates both halves. The 3D pack compiles develop's running-max fix for split-KV (#14276). - **Generator**: - `enumerate.also` sections may use `space: full` and carry manifest `launches`. The enumerator refuses a registry that pairs a reduce when the manifest declares no launches, and keeps one default per cell. - The gates are launch-aware: g1 keys on launches, g3 uses `--subset` for full spaces, g4 compares launch sets, and g6 walks every launch and checks module roles. - g2 and g3 run per engine with the shipped ids reused (#14109), comparing every descriptor byte for byte. - `generate.py --stamp-provenance` stamps `provenance.generated_from` (config, manifest, generator and rocKE commits). It is opt-in, so a default regeneration stays byte-identical. - `variant_reduce` fixes: - it reads the format-2 winner lines the runtime has written since #13028; - it keys a graph on its content with `Graph.id` (minted fresh at each finalization) zeroed, so sweep rounds match. - **Catalogs** (the manifests enumerate the full spaces; the selections are `variant_reduce`'s): - Unified2D: 222 kernels over 198 cells (33 head configs × bf16/fp16 × 3 masks): the narrow nw4/mw16/t64 geometry in every cell, transposed 32x32 nw4/mw32/t64 in 21 and narrow nw2/mw16/t32 in 3. The wide 32x32 geometry won no graph and is not shipped. - Unified3D: 123 kernels: num_segments 16 in all 116 decode cells, 32 or 64 in 7 bf16 bottom-right D128 variants. The K/V cache policy follows #14367's LLC-aware rule, which picks `"all"` for every shipped decode cell. - **Descriptors**: both are generated with develop's generator (post-#13028 UED/UHD, UKD `sources` for 3D), every kernel id reused from the full catalog. A regeneration with the shipped ids reused reproduces them byte for byte. The KDPs live in DVC (`ingestor` remote). - **Shared code** (one home per fact): - `packs/UnifiedAttentionGraphMatch.hpp`: the graph-match helpers both unified packs use (operand/BSHD checks, SDPA node, mask classification, dtype, metadata readers); the tests' graph builder is `UnifiedAttentionTestGraphs.hpp`. - rocKE `kernels/common/attention_mask_type.py` (mask ordinals and band -> kind, used by both tiled specs and re-exported by dispatch) and `attention_split_kv_reduce.py` (the segment -> reduce derivation the dispatcher and the module export share). - CMake registers Dense, Unified2D and Unified3D by one loop per lane (external integration, gpu_ref, census). - **Tests**: - The device knobs suite (`IntegrationGpuGfx950AttentionUnifiedKnobs.cpp`) adds a harness `prepareWorkspace` hook. It forces the default geometry/segment count across the head configurations and each other shipped class on a cell that carries it, and `CoversEveryCatalogKernelClass` holds those tables to the catalog. - The engine-pinned `gpu_ref` targets and external-integration lanes cover both engines. - The ingestor save census names all seven pack dispatch symbols. - The Unified2D/3D integration TOMLs are reserved with the gfx950 census shard (`hkp_reserve_census_shard(... FILES ...)`, as #14573 requires of every shard entry). **Constraints and follow-ups** - #13372 and #13889 are both in develop; this branch no longer carries a copy of #13889. - Heads are catalog cells (baked), window = 0, 2D batch ≤ 511, i32 KV ≤ 0x7FFF0000. D256 is deferred for both engines. The kernarg work is tracked in AICK-2675. - There is no 2D/3D `select_path` crossover replay: the split is the static decode rule above. - Runtimes older than UKD 1.1 skip the 3D descriptors. Strided-KV 3D is a follow-up. - Follow-up: execution-plan save/restore (#13315) for both handlers. Today they report `supportsSaving() == false`, as the RDNA handler does. - One earlier MI355X sweep hit a single GPU fault in Unified2D on `u2d_rocke/rocke_sdpa_035` (bf16, D 128, Hq 64 / Hkv 8, causal, S 512). It has not reproduced since: not in the forced geometries on that cell, 15 single-graph runs, 4 full-corpus runs on MI350X, or the MI355X full-corpus re-check above. - Unified2D kernel time moved with develop's tiled-2D changes (#14542, #14616), not with this PR: at this head it is 1.035× (geomean) its pre-rebase time on the same 188 graphs, up to 1.36× on long causal bf16 D128 prefill (S 8192), while Dense in the same runs is flat (0.997). Dense still ranks first wherever it serves; the table above is at this head. - Measurement note: the pinned dnn-benchmarking (v1 result schema, which `sweep.py` reads) has no bottom-right causal PyTorch reference. The sweeps used upstream dnn-benchmarking #84's reference, with the 9 bottom-right corpus graphs written as `causal_mask: false, right_bound: 0, diagonal_alignment: BOTTOM_RIGHT` (the same band under hipDNN's `resolveDiagonalBand`): #84 reads `causal_mask: true` as top-left regardless of the alignment. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Christopher Millette <63608002+cgmillette@users.noreply.github.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
hip-kernel-provider can now save and load execution plans built by its generic kernel ingestor engines. A saved plan carries the selected kernel's code object (kpack HSACO) and the dispatch handler's launch inputs, so loading a plan skips kernel selection, benchmarking and compilation. This uses hipDNN's existing
to_binary()/to_compiled_plan_binary()save path and its three plugin serialization hooks; hipDNN's plugin API is unchanged. The last commit fixes a data race in the hipDNN backend's serialized-plan cache that this feature makes reachable (ALMIOPEN-2941). Re-saving a plan that was loaded from saved bytes stays refused; it is deferred to ALMIOPEN-2940.JIRA ID : ALMIOPEN-2642
Related: ALMIOPEN-2940 (re-saving a loaded plan; deferred, not in this PR), ALMIOPEN-2941 (backend serialized-plan cache race; fixed in this PR as a separate hipDNN-core commit).
Risk Assessment
Risk 3. New feature path plus a new versioned payload format (FlatBuffers, format 1.0) and new plugin SDK interfaces for dispatch handlers. Behaviour change for existing users: every ingestor engine whose kernels can all be saved now reports
SUPPORTS_EXECUTION_PLAN_SERIALIZATION, soto_binary()of a built plan from such an engine now embeds the plan; a benchmarking plan saved before its first execute is refused instead of saved graph-only. Other engines (HIP_MLOPS, ASM_SDPA, HIP_FLASH2) are unaffected except that a plan-only save from them now returnsHIPDNN_STATUS_PLUGIN_ERROR(NOT_APPLICABLE) rather thanHIPDNN_STATUS_NOT_SUPPORTED, because the hooks are per plugin. Only kpack kernels can be saved (one source-kind gate); restore checks the GPU target, the dispatch handler, the kernel signature and the launch values before loading, and the decoder refuses a payload whose recorded kernel SHA-256 does not match its code object; the gfx950 attention scale range that plan building already enforces is also enforced on this new load path. Save reads code-object bytes through the kpack module cache's existing read path, which is now one function shared by load and save.HIPKERNELPROVIDER_GENERATE_SERIALIZE_HEADERSdefaults ON, except in a TheRock build (detected byTHEROCK_SUBPROJECT_TARGET), where it defaults OFF and the committed generated header is compiled; an explicit-Dalways wins. With the ingestor enabled and the option ON, a configure that cannot findflatc, or finds a separateflatcbinary whose version differs from the expected FlatBuffers version, is an error; setting the option OFF uses the committed generated header. hipDNN core:ExecutionPlanDescriptorgains a mutex, which makes the (internal, not installed) class non-copyable and non-movable; no code copies or moves it.ASIC Coverage
The save/restore code is architecture-independent host code and no kernel, dispatch default or support surface changes, so no multi-arch sweep is required. No PR CI lane exercises this feature on a GPU: the multi-arch GPU lanes build with
HIPDNN_ENABLE_KERNEL_INGESTOR=OFF(the default), and the hipDNN superbuild lanes build it ON on GPU-less runners. GPU verification therefore rests on the manual runs below: the round-trip tests execute real kernels on gfx90a, gfx942 and gfx950 (the gfx950 attention engine's save/load cases only run there). The second-review fixes were tested on gfx942 only; they change no attention code (one comment in the gfx950 attention launch-values header). PR CI shows that the ingestor-disabled build is unaffected and that the feature compiles in the Linux and Windows superbuilds; the Linux superbuild lane also checks that the committed generated header matches its schema.Testing Summary
{nullptr, 0}and{nullptr, 4}; 16 threads loading one restored kernel on a new device load it once (50 repeats); in hipDNN's backend, 16 threads serializing one plan call the plugin once and get identical bytes (50 repeats), and a failed fill is retried.flatcstops with the configure error. Second review: the generation default in a plain and a TheRock-style configure, a standalone provider build against an installed hipDNN with no hipDNN source tree, and a mismatchedflatcversion.gfx942forgfx942:sramecc+:xnack-); archive entries rank an exact key over a generic one; a restored plan is admitted on its base architecture and refused on another device; the serialization note needs every pack's dispatch handler to reportsupportsSaving(), and a test checks each handler's answer against its real save and restore.tryGetEngineNameFromId, and the engine-name helpers that now use it.Testing Checklist
ctest -L "^quick" -E '<3 suites>'- ASIC: gfx942 - Status: Passed (39/39 at e0a3184, MI300X; before the second-review fixes)hip_kernel_provider_tests/hip_kernel_provider_integration_testswith the save/load gtest filters - ASICs: gfx942 (48/48 unit at e0a3184), gfx950 (at 2c325f2) - Status: PassedIntegrationGpuGfx950AttentionDenseSaveLoad.*- ASIC: gfx950 - Status: Passed (2/2 at 2c325f2)-DHIPDNN_ENABLE_KERNEL_INGESTOR=OFF- ASIC: gfx942 - Status: Passed (29/29 at 2c325f2)flatc(configure error), and with generation OFF (committed header used) - Status: Passed (at e0a3184)ENABLE_CLANG_TIDY=ONbuild - gfx90a - Status: Passed (0 diagnostics at 6a7e858)-D: ON;-DTHEROCK_SUBPROJECT_TARGET=hipkernelprovider: OFF; that plus-D…GENERATE_SERIALIZE_HEADERS=ON: ON - Status: Passed (dfed5e4)ROCM_LIBRARIES_ROOTpointing at a missing folder - configures,flatc25.9.23 = expected 25.9.23, header generated with the provider's five flags, byte-identical to the committed one - Status: Passed (dfed5e4)flatcreporting 1.2.3 - configure error naming the path, 1.2.3 and 25.9.23 - Status: Passed (dfed5e4)cmake --build build/debug-gfx942- ASIC: gfx942 - Status: Passed (dfed5e4;flatcran with the five flags, committed header unchanged)hip_kernel_provider_testsper-fix filters - ASIC: gfx942 MI300X - Status: Passed, 0 skipped (33/33 at dfed5e4; 34/34 at b19afb3; 48/48 at 54d4996; 37/37 at 2c5b083; 36/36 at 6d7eeb8; 26/26 +TestPointwiseAddDispatch*10/10 at a2d767c; 5/5 at b194537;TestIngestorKernelCode*27/27 and the new concurrency test 50/50 repeats at 0ea9a72)IntegrationGpuKernelIngestor*(save, load, execute) - ASIC: gfx942 MI300X - Status: Passed, 0 skipped (26/26 at 54d4996, 2c5b083 and a2d767c)DISABLED_PrintLaunchContractSamples- Status: Passed (pointwise, conv and attention samples byte-identical at 2c5b083)hipdnn_backend_tests --gtest_filter="TestExecutionPlanDescriptor.*:TestGraphAndPlanSerializationApi.*", thenSerializeBackendPlanIsSafeFromManyThreadswith--gtest_repeat=50- ASIC: gfx942 MI300X - Status: Passed (61/61 and 50/50, 9dec3ca)-DBUILD_THREAD_SANITIZER=ON,TestExecutionPlanDescriptor.*with--gtest_repeat=20- Status: Passed (39/39 in every repetition, 0 ThreadSanitizer reports, at 3d85e09)flatcversion check, schema-hash pin, kernel-digest check, buffer encoder, full-payload samples, geometry pins, destroy hook, concurrent load, backend cache lock) - ASIC: gfx942 (configure checks and ThreadSanitizer on CPU nodes) - Status: Passed (20/20 caught at 93831e5; restored tree re-tested clean)Check the committed ingestor plan header matches its schema- Status: Passed (PR CI at 3b93962)cmake --build, thenctest -L "^quick" -E '<3 suites>'- ASIC: gfx942 - Status: Passed (65/65 at 6e80ed1)hip_kernel_provider_tests/hip_kernel_provider_integration_testswith the save/load gtest filters - ASIC: gfx942 MI300X - Status: Passed (113/113 unit, 26/26 integration at 91e9fd9)hip_kernel_provider_tests,hip_kernel_provider_integration_tests,hipdnn_plugin_sdk_tests- ASIC: gfx942 MI300X - Status: Passed (51/51 unit, 26/26 integration, 33/33 plugin SDK at b33acfa)hipdnn_plugin_sdk_tests- ASIC: gfx942 MI300X - Status: Passed (784 passed, 0 failed; 2 skipped by environment: 2 GPUs, comma-decimal locale)ENABLE_CLANG_TIDY=ONbuild - gfx90a - Status: Passed (0 diagnostics at 0eeefd0; the fresh run at 9dec3ca found 1, fixed by 0eeefd0)ENABLE_CLANG_TIDY=ONbuild - gfx90a - Status: Passed (0 diagnostics; the code matches b33acfa except one test file)HipDNN Windows Superbuild(ingestor ON, ROCm SDK clang++) - Status: Passed at 6a7e858Technical Changes
kernel_ingestor_engine/serialization/): a 48-byte header (marker, format version, kind, body SHA-256) followed by a FlatBuffers body (ingestor_plan.fbs): code object, source kind, target, kernel signature, dispatch symbol and versioned launch-input values. Every schema field carries an explicitid. Readers accept their own major and any lower minor; damage isINVALID_VALUE, incompatibilityNOT_APPLICABLE. The schema evolution rule (what a minor version may add, what no version may do) is stated iningestor_plan.fbsand RFC 0009, andINGESTOR_PLAN_SCHEMA_SHA256is pinned beside the format version: a test fails on any schema edit that does not also update it (the hash is taken over the LF-normalised schema).HIPKERNELPROVIDER_GENERATE_SERIALIZE_HEADERSON the build regeneratesingestor_plan_generated.hwith the provider's own flatc flags list (serialization/flatc_flags.txt, a copy of hipDNN's list today), so the build no longer readsprojects/hipdnn. A missingflatc, or a separateflatcbinary whose version differs from the expected FlatBuffers version, is a configure error; aflatcbuilt in the same tree is not version-checked. The default is ON, and OFF in a TheRock build. OFF uses the committed header and says so.IngestorPlanCapture): saves the chosen single-kernel plan; refuses a benchmarking plan with no winner, a restored plan, a non-ingestor plan, a non-kpack kernel and a handler that cannot save. The code object is re-read from its kpack archive at save time through the kpack module cache's reader.IngestorPlanRestore,RestoredIngestorPlan): ordered checks (header, body digest, FlatBuffers verifier and structure, kernel SHA-256 against the code object, device target (a saved plan records the base architecture of the GPU it was saved on, also when the kernel came from a generic archive entry, and loads only on that architecture), dispatch handler, handler supports restore, kernel signature, launch-input version and values, module load); the restored plan launches through the handler without selection. The gfx950 attention scale range check that plan building already applies is applied on restore too, through one predicate shared by both paths.kernel_ingestor_engine/packs/): the dispatch-version rule covers the saved values, the kernel arguments the launch passes (order, count, kinds, sizes, meanings) and how the launch derives grid, block and dynamic LDS; it is stated in the pack launch-values headers,SavedDispatch.hpp,PluginDevelopment.mdand RFC 0009. Conv and pointwise compute their launch geometry in header-only helpers (ConvFwdGeometry.hpp,PointwiseGeometry.hpp) called by both build and restore.core/PluginSerialization.cpp): exportshipdnnEnginePluginSerializeExecutionContext,…DestroySerializedExecutionContextand…CreateExecutionContextFromSerializedwhen the ingestor is enabled; every hook runs inside the plugin's exception boundary, and the save hook hands its buffer to hipDNN as its last step, so a failed save leavesserializedContextuntouched. The save encodes straight into thenew[]buffer it hands over (encodeIngestorPlanToBuffer) and frees the payload's code bytes once the builder has them. The destroy hook accepts a nullptr, as the hook contract allows.supportsSaving()), and applied on both of the existing engine-construction paths (the first Container's discovery-built state manager and later Containers).compilation/KpackModuleCache.hpp):load()'s existing read steps (open throughSharedKpackArchives::open(), arch selection, digest check, and eviction on a decompress failure or digest mismatch) move unchanged into onereadVerifiedKpackCodeObject, which save also uses.hipdnn_plugin_sdk/ingestor/,interfaces/IPlan.hpp):IKernelDispatchHandler::saveLaunchInputs/restoreLaunch/supportsSaving(defaults: not supported),SavedDispatch,IPlan::saveablePlan(),BenchmarkPlan::chosenPlan(), descriptor additions.tryGetEngineNameFromId(non-throwing lookup in the engine-name registry). The existinggetEngineNameFromIdandengineNameOrHexnow look up through it, with the same results and exception message;engineNameOrHexno longer throws and catches internally.backend/src/descriptors/ExecutionPlanDescriptor.{hpp,cpp}, ALMIOPEN-2941): one mutex covers the whole check, fill and read of the serialized-plan cache inserializeBackendPlan; the cache is assigned only after the plugin call and the build succeed, so a failed fill is retried on the next call..github/workflows/hipdnn-superbuild-ci.yml): both superbuild configures pass-DHIPKERNELPROVIDER_GENERATE_SERIALIZE_HEADERS=ON; the Linux job then rebuilds the header target and runsgit diff --exit-codeon the committed header.PluginDevelopment.mdand the provider README describe saving, including that a loaded plan runs the code it was saved with (to use a kernel fixed by an upgrade, build the plan again from its graph and save it) and that a save may fail when the kpack archive changes after the plan build.projects/hipdnn/.clang-tidy, plugin SDK headers,projects/hipdnn/docsand the backend'sExecutionPlanDescriptorare touched, so hipDNN-core code owners are needed.Commits are kept per step for review: payload format, its build option and schema, capture, restore, hooks, engine-identity removal, tests and docs, test consolidation, the scale check on restore with the fast-path note test, the shared kpack read path, and the review fixes (save-buffer hand-over order, capability-probe test, explicit schema field ids, configure error on a skipped regeneration, engine-name lookup, resealed-byte sweep, launch-contract samples, comment corrections). The second review adds one commit per item: header-generation flags, default and version check (dfed5e4); schema-hash pin (b19afb3, rule wording 93831e5); kernel-digest check (54d4996, formatting 041ceca); encode into the final buffer (2c5b083); full-payload samples (6d7eeb8); geometry pins and the dispatch-version rule (a2d767c); archive wording (e17d7a0); destroy hook (b194537, clang-tidy const fix 0eeefd0); concurrent-load test (0ea9a72); and, last and self-contained so it can be split out, the hipDNN backend cache lock (9dec3ca, ALMIOPEN-2941). The branch is rebased onto develop
dd76b08821c.🤖 Generated with Claude Code