Repository navigation
feat(miopen-provider): support standalone binary pointwise - #11934
Merged
averinevg merged 24 commits intoSep 22, 2026
Merged
Conversation
Adds ADD/SUB/MUL/MAX_OP/MIN_OP as standalone binary pointwise ops backed by miopenOpTensor, with applicability checks, plan/builder, and shared pointwise tensor-checking helpers reused from the unary activation path. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Adds IntegrationGpuPointwiseBinary GPU test suite covering elementwise add/sub/mul/max/min across 3d/4d/5d and broadcast shapes for fp32/fp16. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…ary-elementwise-op-support
…ary-elementwise-op-support # Please enter a commit message to explain why this merge is necessary, # especially if it merges an updated upstream into a topic branch. # # Lines starting with '#' will be ignored, and an empty message aborts # the commit.
…ary-elementwise-op-support
…ulti-node test IsSupportedFalseForMultiNodeGraph only stubbed nodeCount(), so it stayed green even with the nodeCount != 1 guard removed: the next call, hasOnlySupportedAttributes(), was left unstubbed and gmock silently returned its default value instead of failing. Assert it is never reached instead, which pins the guard's short-circuit behavior. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…nit tests Merge the graph/execution helper headers into a single PointwiseCommon.hpp, drop unit tests that only duplicated coverage already provided end-to-end by the integration suite (shape/rank sweeps, per-mode numeric execution), and consolidate ad hoc fixtures. Kept and added the tests that exercise behavior unique to the unit level: defensive internal-error paths and direct plan construction from a valid graph. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Replace the miopen-provider-local PointwiseGraphSpec/createPointwiseGraph with createValidPointwiseGraph in hipdnn_test_sdk's FlatbufferGraphTestUtils.hpp, following the SDK's plain-parameter convention instead of a config struct. This fills in a stale placeholder that was waiting for exactly this helper, and lets other providers reuse it instead of duplicating graph-building logic. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…e checks Drop the numbered check-index comments and ones that only restated the following line; keep the comments explaining non-obvious rationale (load-bearing invariants, why a check exists, design decisions). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…ary-elementwise-op-support
Redesign createValidPointwiseGraph (renamed createPointwiseGraph) into a fully generic SDK helper with per-tensor dims/strides grouping, full third-input support, and the remaining activation params. Add createUnaryPointwiseGraph/createBinaryPointwiseGraph wrappers that expose only the relevant subset of params, and migrate call sites to use them. Eliminate the now-trivial validBinaryGraph pass-through in PointwiseCommon.hpp, trim unused includes/usings there, and rename the Max/Min mode-case labels. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
2 of 5 tasks
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## develop #11934 +/- ##
===========================================
- Coverage 70.75% 70.35% -0.40%
===========================================
Files 2849 2815 -34
Lines 468199 463759 -4440
Branches 69164 68269 -895
===========================================
- Hits 331243 326269 -4974
- Misses 112960 113873 +913
+ Partials 23996 23617 -379
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
adickin-amd
reviewed
Sep 10, 2026
adickin-amd
reviewed
Sep 10, 2026
adickin-amd
reviewed
Sep 10, 2026
adickin-amd
reviewed
Sep 10, 2026
adickin-amd
reviewed
Sep 10, 2026
adickin-amd
reviewed
Sep 10, 2026
adickin-amd
reviewed
Sep 10, 2026
6 tasks done
…ary-elementwise-op-support
…ary-elementwise-op-support
Adds a graph-only quick/Pointwise/Binary bundle (40 cases across ADD/SUB/MUL/MAX_OP/MIN_OP, fp16/fp32, and multiple ranks/layouts), captured from IntegrationGpuPointwiseBinary and placed via place_bundles.py, so CI exercises the binary pointwise path without requiring -DBUILD_CPP_GRAPH_TESTS=ON. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
… in the deterministic engine miopenActivationForward and miopenOpTensor have no algorithm selection, atomics, or cross-thread reduction reachable, so both are deterministic unconditionally and need no deterministic-mode flag (unlike conv). Adds GPU fixtures asserting bit-exact reproducibility for both, mirroring the existing conv determinism tests, and updates the shared EngineOrdering comments and OperationSupport.md to match. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…gainst int64 overflow Dims are validated positive but not bounded, so the numel and packed-stride running products could overflow int64_t before the INT32_MAX guard saw them -- undefined behaviour, and a wrapped value can land back under the limit and pass a check that is load-bearing for MIOpen's int-narrowed work_per_wg. Check before each multiply instead. Also covers checkPacked, which runs on in_0/in_1 before their dims are compared against the output's. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
buildPlan was called by no unit test: the builder suite covered isApplicable, getMaxWorkspaceSize and getCustomKnobs, and the plan suite constructed MiopenBinaryPointwisePlan directly, so nothing asserted the builder hands a constructed plan to the execution context. Assert hasValidPlan() and the plan's zero workspace size, per the dominant pattern across the conv and batchnorm builder suites. Lives in the mode-parameterized fixture, so all five modes are covered. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
createPointwiseGraph had grown to 24 defaulted parameters, with callers padding long runs of std::nullopt and hand-writing /*paramName=*/ labels to stay readable. Replace the parameter list with a PointwiseGraphSpec whose field order and defaults mirror the old signature exactly, and fold the two positional wrappers into PointwiseGraphSpec::unary()/::binary(), which keep their previous parameter lists. Callers that need an operand neither helper exposes now set the named field instead. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…rip fixture The pointwise arm of TestJson.GraphToJsonAndBack is the only place JSON ever serializes a PointwiseAttributes node, so axis_tensor_uid, in_1/in_2, the relu/swish/elu/softplus knobs and a non-FLOAT tensor dtype are round-tripped nowhere else. Repointing it at the valid-by-default builder left all of those null or default on both sides of the trip, where they cannot catch a writer/reader bug. Add PointwiseGraphSpec::fullyPopulated() to set them distinctly again, plus the axisTensorUid field the builder previously hardcoded to nullopt. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…ary-elementwise-op-support
adickin-amd
approved these changes
Sep 21, 2026
adickin-amd
left a comment
adickin-amd
left a comment
Contributor
There was a problem hiding this comment.
LGTM, thanks for addressing my concerns.
Contributor
TheRock Submodule Bump ActivityNewest first
|
evedovelli
pushed a commit
that referenced
this pull request
Sep 23, 2026
## Summary Adds ADD, SUB, MUL, MAX_OP, MIN_OP as standalone single-node pointwise graphs to the hipDNN MIOpen provider, via `miopenOpTensor`. Frontend, schema, and CPU reference already supported these modes; this closes the provider-side gap. ISSUE ID : #8458 ## Risk Assessment Level 3 — new feature path, additive. The one behavior-adjacent change is covered by the unchanged unary suite passing. ## ASIC Coverage Standard PR CI is sufficient — arithmetic-only op with no arch-specific dispatch. ## Testing Summary - New unit tests for the resolver, plan, and plan builder (`TestMiopenBinaryPointwise*.cpp`) and the shared tensor validator (`TestMiopenPointwiseTensorChecks.cpp`). - Unary suite extended to assert it rejects two/three-input nodes and binary modes. - New GPU smoke test (`IntegrationGpuPointwiseBinary.cpp`) over a shape × mode matrix. - Test-SDK helper (`FlatbufferGraphTestUtils.hpp`) generalized for two-input graphs. ## Testing Checklist - [x] `miopen_plugin_tests` - [x] `miopen_plugin_integration_tests` - [x] PR CI ## Technical Changes - Mode mapping: ADD/MAX_OP/MIN_OP map directly; SUB is Add with the second operand scaled by −1. Beta fixed at 0. - Resolver (`MiopenBinaryPointwiseChecks.{hpp,cpp}`): single node, rank 3–5, matching ranks, packed channels-first strides, only the second operand may broadcast, no in-place, FP16/FP32 only, numel ≤ INT32_MAX. - Shared tensor validator (`MiopenPointwiseTensorChecks.{hpp,cpp}`, new): dtype/null/shape checks factored out of the unary path and reused here. - Plan + builder (`MiopenBinaryPointwisePlan*`, `MiopenBinaryPointwisePlanBuilder*`): zero workspace, registered in `MiopenContainer.cpp`, appended last. - Unary builder now rejects nodes with a second/third input (previously silently dropped it). - Docs: both `OperationSupport.md` files and `miopen-plugin.rst` updated with the new row. --------- Co-authored-by: Claude Sonnet 5 <noreply@anthropic.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
Adds ADD, SUB, MUL, MAX_OP, MIN_OP as standalone single-node pointwise graphs to the hipDNN MIOpen provider, via
miopenOpTensor. Frontend, schema, and CPU reference already supported these modes; this closes the provider-side gap.ISSUE ID : #8458
Risk Assessment
Level 3 — new feature path, additive. The one behavior-adjacent change is covered by the unchanged unary suite passing.
ASIC Coverage
Standard PR CI is sufficient — arithmetic-only op with no arch-specific dispatch.
Testing Summary
TestMiopenBinaryPointwise*.cpp) and the shared tensor validator (TestMiopenPointwiseTensorChecks.cpp).IntegrationGpuPointwiseBinary.cpp) over a shape × mode matrix.FlatbufferGraphTestUtils.hpp) generalized for two-input graphs.Testing Checklist
miopen_plugin_testsmiopen_plugin_integration_testsTechnical Changes
MiopenBinaryPointwiseChecks.{hpp,cpp}): single node, rank 3–5, matching ranks, packed channels-first strides, only the second operand may broadcast, no in-place, FP16/FP32 only, numel ≤ INT32_MAX.MiopenPointwiseTensorChecks.{hpp,cpp}, new): dtype/null/shape checks factored out of the unary path and reused here.MiopenBinaryPointwisePlan*,MiopenBinaryPointwisePlanBuilder*): zero workspace, registered inMiopenContainer.cpp, appended last.OperationSupport.mdfiles andmiopen-plugin.rstupdated with the new row.