feat(hipdnn): implement RFC 0016 runtime pass-by-value tensors#9286
Draft
SamuelReeder wants to merge 3 commits into
Draft
feat(hipdnn): implement RFC 0016 runtime pass-by-value tensors#9286SamuelReeder wants to merge 3 commits into
SamuelReeder wants to merge 3 commits into
Conversation
Implements the single-flag 2-bit model from RFC 0016: per-tensor is_runtime_pass_by_value schema flag, backend attributes/descriptor support, wrapper accessors, cuDNN-parity frontend API (ScalarType, constructors, setters, getters), execute-time variant-pack filter, plugin API version floor (1.2.0) + filter, and serialized-graph reader-version guard. JIRA ID : ALMIOPEN-2240
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
SamuelReeder
commented
Jul 10, 2026
Comment on lines
+95
to
+98
| * Matches cuDNN's plain scalar constructor: the tensor is runtime | ||
| * pass-by-value with a baked default (runtime flag set), flooring the | ||
| * provider at 1.2.0. Use ScalarType::COMPILE_TIME_CONST or | ||
| * set_compile_time_constant() for a baseline-1.0.0 baked constant. |
Contributor
Author
There was a problem hiding this comment.
I don't think we need to be this descriptive. The rfc already outlines the motivation.
SamuelReeder
commented
Jul 10, 2026
Comment on lines
+104
to
+111
| for(const auto& tensor : graph->tensors) | ||
| { | ||
| if(tensor && tensor->is_runtime_pass_by_value) | ||
| { | ||
| anyRuntimePassByValue = true; | ||
| break; | ||
| } | ||
| } |
Contributor
Author
There was a problem hiding this comment.
Can extract the helper for this and share with the engine plugin resource manager.
…6 runtime pass-by-value scalars
Extends both in-tree providers to keep serving existing graphs and support
a newly runtime user-supplied host scalar (epsilon/momentum) after the
core RFC 0016 breaking change (ALMIOPEN-2299, ALMIOPEN-2238).
Phase A - version bumps:
- MIOpen: HIPDNN_PLUGIN_API_VERSION 1.0.0 -> 1.2.0.
- hip-kernel: define HIPDNN_PLUGIN_API_VERSION "1.2.0" (was previously
unset, defaulting the host to 1.0.0).
Phase B - execute-time host-scalar read, both providers:
- New ScalarOperand/makeScalarOperand/resolveScalarOperand helper trio in
MiopenUtils.{hpp,cpp} and hip-kernel's core/Utils.{hpp,cpp}. Pure
runtime-user-supplied scalars (is_runtime_pass_by_value() &&
value_type()==NONE) resolve from the execute-time device_buffers slot;
every other state (compile-time constant, runtime-with-default) keeps
using the baked value and ignores any device_buffers slot for that uid.
- Converted every batchnorm/layernorm/RMSnorm plan that reads epsilon or
momentum to build a ScalarOperand at plan-build and resolve it at
execute, in both providers.
- Fixed a pre-existing applicability-check bug (both providers, batchnorm/
layernorm/RMSnorm): the "skip scalar tensors" filter in
checkTensorLayoutsAndDimsSupported checked only value_type()==NONE,
which incorrectly treated a pure runtime-user-supplied scalar (also
value_type()==NONE, by design) as a real tensor and rejected it on the
3D/4D/5D dimension check. Filter now also skips
is_runtime_pass_by_value() tensors.
Phase C - tests:
- hip-kernel: added ApiVersion unit tests (TestPluginPublic.cpp) mirroring
MIOpen's existing coverage; MIOpen's existing test already covers the
1.2.0 bump.
- New GPU integration regression test per provider
(IntegrationGpuBatchnormPassByValueEpsilon.cpp): proves a runtime host
scalar delivered via set_as_runtime_parameter() at execute produces the
same output as an equal-valued compile-time constant, and that a
different runtime value visibly changes the output.
Verified: both providers build clean under the superbuild presets
(hipdnn-providers, hip-kernel-provider); full unit suites pass
(1048/1048 MIOpen, 368/368 hip-kernel); full GPU integration suites pass
(125/125 MIOpen non-skip, 1862/1862 hip-kernel non-skip) on gfx90a
(MI210), including the new PassByValueEpsilon regression tests.
SamuelReeder
commented
Jul 10, 2026
Contributor
Author
There was a problem hiding this comment.
Confirm compile time is tested too.
…red RuntimePassByValue.hpp Consolidates provider-duplicated RFC 0016 runtime pass-by-value helpers into plugin_sdk, the natural home for cross-provider plugin-side infrastructure (plugin_sdk already hosts PluginHelpers.hpp-style generic helpers and is already linked against hipdnn_flatbuffers_sdk + hipdnn_data_sdk). - New header-only hipdnn_plugin_sdk/RuntimePassByValue.hpp: findDeviceBuffer (previously copy-pasted verbatim in MIOpen, hip-kernel, and hipblaslt), plus ScalarOperand/makeScalarOperand/resolveScalarOperand (previously duplicated in MIOpen and hip-kernel). - MiopenUtils.hpp, hip-kernel core/Utils.hpp, and HipblasltUtils.hpp now `using`-import the shared symbols into their provider namespaces so all existing call sites (miopen_utils::findDeviceBuffer, etc.) keep working unchanged. Removed the now-dead duplicated .cpp implementations. - Added dedicated plugin_sdk unit tests (TestRuntimePassByValue.cpp) covering findDeviceBuffer edge cases and all three ScalarOperand states (compile-time constant, runtime-with-default, pure runtime user-supplied) including every resolveScalarOperand dtype branch and error paths. Removed the resulting duplicate findDeviceBuffer tests from TestMiopenUtils.cpp, hip-kernel's TestUtils.cpp, and TestHipblasltUtils.cpp. Rationale: two independent providers had already drifted on interpreting the RFC 0016 by-value contract (the checkTensorLayoutsAndDimsSupported scalar-skip bug fixed in the prior commit was found and had to be fixed separately in both MIOpen and hip-kernel) -- centralizing the contract logic removes that class of bug for any current or future provider. Verified: both providers + hipblaslt build clean under the superbuild presets. Full unit suites pass (1046/1046 MIOpen, 363/363 hip-kernel, 125/125 hipblaslt, 68/68 plugin_sdk incl. 13 new RuntimePassByValue tests -- counts down from the prior commit only by the removed duplicate findDeviceBuffer tests). Full GPU integration suites pass unchanged (125/125 MIOpen, 1862/1862 hip-kernel, 335/335 hipblaslt non-skip) on gfx90a (MI210), including both PassByValueEpsilon regression tests.
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
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Implements RFC 0016 (single-flag 2-bit model) in hipDNN core: a per-tensor
is_runtime_pass_by_valueschema flag, backend attribute/descriptor support, wrapper accessors, the cuDNN-parity frontendTensorAttributesAPI (ScalarType, constructors, setters, getters), the execute-time variant-pack filter, a plugin API version floor (1.2.0) with filtering, and a serialized-graph reader-version guard. Followed by PR #8860 (RFC merge).JIRA ID : ALMIOPEN-2240
Risk Assessment
High risk.
get_pass_by_value()changes return shape and semantics (value now surfaces only for the "runtime with default" state, not for compile-time constants), which is a breaking change to a public frontend API touching every in-tree caller. The flatbuffer schema changes are append-only defaulted fields (wire-compatible), and the change is fully exercised by 4453 passing local tests across unit, integration, backend, and flatbuffer suites, but the breadth of the getter semantic flip and its cascading effect on execution-test epsilon/scalar construction leaves residual risk for out-of-tree callers not covered by this repo's test suite.ASIC Coverage
ASIC-independent. This is frontend/backend descriptor, schema, and plugin-version-filter wiring — it does not touch kernel selection, op support surface, or provider dispatch behavior for any specific architecture. No provider ops are added or changed. Standard PR CI is sufficient; no dedicated multi-arch sweep is required.
Testing Summary
TensorAttributes,TensorValueAttributes,DescriptorHelpers,DescriptorUnpackHelpers,PassByValueExecuteFilter) covering all three by-value states, the getter/setter matrix, validation rejections, and pack/unpack round-trips.<1.2.0plugins reject a runtime-pbv graph; positive: a1.2.0fake plugin serves it) and the reader-version guard, plus full regression across Batchnorm/Layernorm/RMSNorm/SDPA lifting and lowering suites.TensorDescriptorattribute round-trips (1306/1307/1308) andGraphDescriptormin_reader_versionstamping/rejection.isRuntimePassByValue()/isByValue()/hasCompileTimeConstant()across all flag×value quadrants.Testing Checklist
./build/release/bin/hipdnn_frontend_tests- Status: Passed (1889 tests)./build/release/bin/hipdnn_public_frontend_tests- Status: Passed (355 tests)./build/release/bin/hipdnn_public_backend_tests- Status: Passed (189 tests)./build/release/bin/hipdnn_flatbuffers_sdk_tests- Status: Passed (132 tests)./build/release/bin/hipdnn_test_sdk_tests- Status: Passed (1771 tests, 3 pre-existing skips)Technical Changes
is_runtime_pass_by_value(defaulted, wire-compatible) toTensorAttributesintensor_attributes.fbs, andmin_reader_versiontograph.fbs.HIPDNN_ATTR_TENSOR_CONSTANT_VALUE(alias ofVALUE_EXT) andHIPDNN_ATTR_TENSOR_IS_RUNTIME_PASS_BY_VALUE(1308), with get/set support inTensorDescriptorand string-name lookup.isRuntimePassByValue(),isByValue()(umbrella), andhasCompileTimeConstant()to theTensorAttributesWrapperinterface and implementation.frontend::TensorAttributesto the three-state model (runtime user-supplied, runtime-with-default, compile-time constant): newScalarTypeenum, two-argument constructor,set_compile_time_constant/set_as_runtime_parameter/set_is_pass_by_valuesetters, and state-gatedget_pass_by_value/get_compile_time_constant/get_is_pass_by_value/get_is_runtime_pass_by_valuegetters; tightensvalidate()to reject virtual+runtime-flag and virtual+value independently.Graph::tensor(scalar, ScalarType)factory overload and an execute-timedetail::validatePassByValueVariantPackfilter wired into the fresh-build path ofGraph::execute(), rejecting a variant-pack override of a baked value and requiring a host-supplied entry for a pure runtime scalar.K_PASS_BY_VALUE_MIN_API_VERSION = "1.2.0"and bumpsHIPDNN_ENGINE_API_VERSION_MINORto 2;EnginePluginResourceManager::computeMinimumPluginApiVersionnow floors at 1.2.0 whenever any tensor in the finalized graph has the runtime flag set, andgetApplicableEngineIdsfilters plugins accordingly.min_reader_versionstamp on graph serialization (1 when any tensor is runtime pass-by-value, else 0) and a reject guard on deserialization when the serialized graph requires a newer reader version than the build supports.get_pass_by_value()API (lifting/lowering tests,makeTensorAttributeshelper,FrontendGraphFactory, execution-test epsilon construction) to the correct new getter for their actual semantic intent (compile-time constant vs. runtime-with-default).