Skip to content

feat(miopen-provider): support standalone resample forward (pooling) - #14984

Open
averinevg wants to merge 15 commits into
developfrom
users/evaverin/ALMIOPEN-2684/miopen-provider-route-pooling-fwd
Open

averinevg wants to merge 15 commits into
developfrom
users/evaverin/ALMIOPEN-2684/miopen-provider-route-pooling-fwd

Conversation

@averinevg

@averinevg averinevg commented Nov 3, 2026 •

Copy link
Copy Markdown
Contributor

Routes single-node ResampleFwd graphs (max pooling and both average-pooling modes) on MIOPEN_ENGINE and MIOPEN_ENGINE_DETERMINISTIC to miopenPoolingForward. Max pooling can also write hipDNN's index tensor, and the index rules set here are the contract pooling backward will consume. Graphs MIOpen cannot run as hipDNN defines them are declined rather than approximated.

JIRA ID : ALMIOPEN-2684

Risk Assessment

Level 3. This is a new feature path in a provider: a new plan builder registered on both MIOpen engines, behind applicability checks that decline anything outside the tested set.

ASIC Coverage

The provider code is host-only: graph checks, descriptor setup and calls into MIOpen. Kernel behavior on each GPU architecture is MIOpen's responsibility, so the provider supports ResampleFwd wherever MIOpen does. That means PR CI and the local gfx942 run cover the provider's logic:

  • PR CI runs the standard tier on gfx94X (Linux) and gfx110X (Windows).
  • The local gfx942 run executed the provider suites and the MIOPEN_ENGINE bundle lane.

The new support claims name gfx90a, gfx942 and gfx950 on Linux. The integration lanes enforce them wherever they run. No Windows bundle claims are added.

Testing Summary

  • Test SDK (hipdnn_test_sdk_tests): ResampleFwdGraphSpec and createResampleFwdGraph(), which build the configurable single-node ResampleFwd graphs the provider tests use.
  • Provider unit tests (miopen_plugin_tests) keep only the tests that fail when the provider's own code is wrong; how MIOpen pools is MIOpen's to test:
    • the pooling descriptor wrapper;
    • applicability: one test per decline rule, matching its message, plus the boundaries and transformations the provider computes itself (padding against the window, the padding-mode rule, the INT8 index range, the channels-last tile limit, stride normalization, canonical strides, the layout choice and the resolved configuration);
    • the plan: exact-value runs that catch a wrong descriptor, data type, strides, alpha or beta, index copy or workspace, plus the execution guards;
    • the plan builder: the override-shape decline, deferring to the checks, the workspace size it reports and its errors;
    • both MIOpen engines accepting a ResampleFwd graph.
  • Bundles (hipdnn_integration_tests):
    • a new quick bundle, ResampleFwd/SymmetricPadding (47 cases: all three modes, NCHW/NHWC/NCDHW/NDHWC, fp32/fp16/bf16), because most ResampleFwd/Default cases use asymmetric padding, which MIOpen declines;
    • MIOPEN_ENGINE claims on the zero-padding max-pool cases of the existing ResampleFwd/Default bundles in every tier, which write an index;
    • in auto mode they are verified against the GPU resample reference. The harness compares Y only, because it reads a ResampleFwd index as an input; the index values are checked by the provider's plan tests.

Testing Checklist

  • Test SDK tests - hipdnn_test_sdk_tests - Status: Passed (2349 passed, 4 skips that predate the change)
  • MIOpen provider unit tests - miopen_plugin_tests --gtest_filter='*MiopenPoolingDescriptor*:*MiopenResampleFwd*:TestGpuMiopenContainer.*' - Devices: gfx942 - Status: Passed (104/104)
  • ResampleFwd bundles on MIOPEN_ENGINE, auto mode - hipdnn_integration_tests --test-article <libmiopen_plugin.so> --test-engine MIOPEN_ENGINE --test-config dnn-providers/miopen-provider/config/MIOPEN_ENGINE.toml --gtest_filter='*ResampleFwd*' - Devices: gfx942 - Status: Passed (56/56 claimed cases verified by the GPU reference; the 39 unclaimed cases are declined and skipped)
  • Support overview - python3 dnn-providers/integration-tests/scripts/support_features/feature_matrix.py --check - Status: Passed (overview regenerated and current)
  • Support-matrix tests on the real bundle tree, as hipdnn-superbuild-ci runs them - python3 -m pytest -q dnn-providers/integration-tests/scripts/test_render_support_matrix.py dnn-providers/integration-tests/scripts/test_support_features.py - Status: Passed (1987 passed)
  • pre-commit - pre-commit run --files <changed files> - Status: Passed (the repo's pre-commit config excludes the bundle tree, so verify_support_claims.py and render_support_matrix.py --output - were also run directly: Passed)
  • PR CI - GitHub PR checks - Status: Pending

Technical Changes

  • Pooling descriptor. MiopenPoolingDescriptor owns an N-d miopenPoolingDescriptor_t. It sets the index type and the Image index mode after miopenSetNdPoolingDescriptor, because that call resets both. Image mode records a flattened spatial offset into the unpadded input, which is hipDNN's convention. MIOpen defaults to Mask mode, which its Nd solvers reject for max pooling.
  • Pooling forward plan. MiopenResampleFwdPlan calls miopenPoolingForward with alpha 1 and beta 0.
    • Its workspace is sized by miopenPoolingGetWorkSpaceSizeV2, which includes the transpose buffers for channels-last tensors. Passing the index tensor as the workspace would be undersized under NHWC.
    • When an index is requested, the plan copies the first yElements * sizeof(index) bytes of the workspace into the index tensor on MIOpen's stream.
  • Applicability checks and plan builder. MiopenResampleFwdChecks and MiopenResampleFwdPlanBuilder decline:
    • NOT_SET mode;
    • asymmetric padding, or pre_padding >= window;
    • rank other than 4 or 5;
    • datatypes other than fp32, fp16 and bf16, a non-FP32 compute type, or mixed X/Y types or layouts;
    • non-packed tensors, or a Y that is not the floor-mode output size;
    • multi-node and override-shape graphs;
    • an index on channels-last tensors: MIOpen's index buffer stays packed NC(D)HW while its transposing solvers move only X and Y;
    • an index tensor without generate_index, an index tensor with average pooling, or generate_index without an index tensor for max pooling;
    • an INT8 index when the input spatial size exceeds 128;
    • channels-last tiles beyond the scratch limit of MIOpen's tiled kernels, which have no naive fallback for those layouts.
  • generate_index and average pooling. Average pooling with generate_index set and no index tensor is accepted, and the flag is ignored. The frontend builds that graph, and average pooling produces no index.
  • Padding mode. MAXPOOL with NEG_INF_PAD and average pooling with ZERO_PAD are accepted with any padding. Any other combination, PADDING_NOT_SET included, is accepted only when all padding is zero, because the mode then has no effect on the result. This deliberately accepts ZERO_PAD max pooling and NEG_INF_PAD average pooling when there is no padding.
  • Engine registration. The builder is registered on both MIOPEN_ENGINE and MIOPEN_ENGINE_DETERMINISTIC. MIOpen's pooling forward kernels compute each output in one thread, without atomics.
  • Test SDK. ResampleFwdTensorSpec, ResampleFwdGraphSpec (packed(), makeChannelsLast(), setPadding()) and createResampleFwdGraph() join FlatbufferGraphTestUtils.hpp, with tests in TestFlatbufferGraphTestUtils.cpp.
  • Docs. The MIOpen provider OperationSupport.md, hipDNN OperationSupport.md, miopen-plugin.rst and plugins.rst list the op and its constraints, and note that max pooling ignores NaN inputs. SUPPORT_MATRIX_OVERVIEW.md is regenerated.

averinevg and others added 10 commits November 2, 2026 15:10
MiopenPoolingDescriptor owns an N-d miopenPoolingDescriptor_t for 2 or 3
spatial dimensions. The index type and the workspace index mode are set
after miopenSetNdPoolingDescriptor, which resets both. Indices are always
recorded in Image mode, a flattened spatial offset into the unpadded
input, which is hipDNN's index convention. If a setter fails after the
descriptor is created, the descriptor is destroyed before rethrowing.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Add ResampleFwdTensorSpec, ResampleFwdGraphSpec and createResampleFwdGraph()
to the test SDK (FlatbufferGraphTestUtils.hpp), next to
createValidResampleFwdGraph(). They describe a configurable single-node
ResampleFwd graph. The defaults are a valid fp32 max pool; packed() builds
a pooling graph with a floor-mode Y and the padding mode that matches the
resample mode, optionally with an INT8 or INT32 index; makeChannelsLast()
gives X and Y channels-last strides. Each tensor can be made virtual,
pass-by-value, ragged or strideless, or left out of the graph, so provider
tests can build the graphs their checks decline.

createValidResampleFwdGraph() stays; moving its callers to the spec is
left for a follow-up.

TestFlatbufferGraphTestUtils.cpp covers packed(), makeChannelsLast() and
each branch of createResampleFwdGraph().

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
resample_fwd_applicability::resolve() maps a single-node ResampleFwd graph
to the MIOpen pooling configuration that the plan executes. If the graph is
not supported it throws BAD_PARAM naming the reason; isSupported() logs
that reason and declines.

Accepted graphs:
- MAXPOOL, AVGPOOL_EXCLUDE_PADDING or AVGPOOL_INCLUDE_PADDING, with FP32
  compute;
- FLOAT, HALF or BFLOAT16 X and Y of one type, rank 4 or 5;
- symmetric padding smaller than the window, and a floor-mode Y;
- NEG_INF_PAD for MAXPOOL and ZERO_PAD for average pooling; any other
  padding mode only when all padding is zero;
- an INT8 or INT32 image-mode index, for max pooling in NCHW/NCDHW only.
  Average pooling with generate_index set but no index tensor runs without
  one: average pooling produces no index.

X and Y are handed to MIOpen with the canonical packed strides of the
layout they share. MIOpen infers each tensor's layout from its strides and
picks native or transposing kernels per tensor, but sizes the transpose
workspace from Y alone, so the two must never be classified differently.

A spatial dimension with a single output gets the smallest stride that
still gives one output. This changes nothing numerically, but keeps
MIOpen's workspace and tile estimates sane for global pooling.

Channels-last graphs are declined when their per-lane tile exceeds the
scratch limit of MIOpen's tiled kernels, which have no naive fallback for
these layouts.

The unit tests check each rule once, matching its decline message, plus
the boundaries and transformations the checks compute themselves: padding
against the window, the padding-mode rule, the INT8 index range, the
channels-last tile limit, stride normalization, canonical strides and the
layout choice. They build their graphs with the test SDK's
ResampleFwdGraphSpec.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ngForward

MiopenResampleFwdPlan runs a resolved ResampleFwd configuration with
miopenPoolingForward, using alpha 1 and beta 0. Its workspace is MIOpen's
own, sized by miopenPoolingGetWorkSpaceSizeV2. That covers the indices
and, for channels-last tensors, the transposed copies of X and Y.

When the node generates an index, MIOpen leaves it in the first
yElements * sizeof(index) bytes of the workspace, packed in NC(D)HW order.
The plan copies it into the index tensor on MIOpen's stream. MIOpen never
checks the size of the workspace it is given and knows nothing of the
index tensor, so the plan throws INVALID_VALUE when a required workspace
or the index buffer is missing.

The tests run the plan on the GPU with values worked out by hand. Each
fails when the plan's own code is wrong: the descriptor and tensors it
builds, the workspace it reports and checks, alpha and beta, the index it
copies out, or the canonical strides it hands MIOpen.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…IOpen engines

MiopenResampleFwdPlanBuilder routes a standalone ResampleFwd node to
MiopenResampleFwdPlan. isApplicable declines graphs with override shapes
and otherwise defers to the ResampleFwd checks. getMaxWorkspaceSize and
buildPlan resolve the configuration through the same checks, so the
workspace the builder reports is the one the plan asks for, and an
unsupported graph throws BAD_PARAM from both. The builder has no
execution settings and no custom knobs.

The builder is registered on MIOPEN_ENGINE and MIOPEN_ENGINE_DETERMINISTIC.
MIOpen's pooling forward kernels compute each output in one thread
without atomics, so their results do not vary from run to run.

The tests cover the override-shape decline, the deferral to the checks,
a channels-last max-pooling build whose plan workspace equals the
reported maximum, the BAD_PARAM errors, and both engines accepting a
ResampleFwd graph from the real container.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The MIOpen provider's OperationSupport.md, hipDNN's OperationSupport.md,
miopen-plugin.rst and plugins.rst list standalone resample forward on the
MIOpen provider. The provider notes give what the applicability checks
accept: the modes, ranks, layouts and datatypes, the padding rules, the
floor-mode output size and the index rules. The deterministic engine
section now includes resample forward.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
MIOPEN_ENGINE can run few of the existing ResampleFwd cases: Variant2
uses asymmetric padding, and Default is max pooling with an index. The
new quick bundle ResampleFwd/SymmetricPadding adds 47 cases covering all
three modes, four layouts and three datatypes.

MIOPEN_ENGINE claims those cases and the Default cases it runs in every
tier, on gfx90a, gfx942 and gfx950 (Linux). One quick Default case joins
ffm-quick. SUPPORT_MATRIX_OVERVIEW.md is regenerated.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…IOPEN-2684/miopen-provider-route-pooling-fwd

# Conflicts:
#	dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md
#	projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp
#	projects/hipdnn/test_sdk/tests/CMakeLists.txt

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
flatbuffers::Optional is std::optional here, so the toFlatbuffers lambda
unwrapped and rewrapped the same type. clang-tidy rejects that
(bugprone-optional-value-conversion) and fails the build.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
txt

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…IOPEN-2684/miopen-provider-route-pooling-fwd

# Conflicts:
#	dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@therock-pr-bot

therock-pr-bot Bot commented Nov 3, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

therock-pr-bot Bot commented Nov 3, 2026

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

averinevg and others added 2 commits November 3, 2026 19:34
…test

clang-tidy misc-unused-using-decls flagged it as an error in CI.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…IOPEN-2684/miopen-provider-route-pooling-fwd
@codecov

codecov Bot commented Nov 4, 2026 •

Copy link
Copy Markdown

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files
@@             Coverage Diff             @@
##           develop   #14984      +/-   ##
===========================================
- Coverage    78.09%   77.69%   -0.40%     
===========================================
  Files         3005     2905     -100     
  Lines       521953   490948   -31005     
  Branches     81452    74312    -7140     
===========================================
- Hits        407577   381419   -26158     
+ Misses       85458    84566     -892     
+ Partials     28918    24963    -3955     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 51.04% <ø> (-2.85%) ⬇️ Carriedforward from df16ab4
TensileLite-Unit 84.35% <ø> (-0.72%) ⬇️ Carriedforward from df16ab4
hipBLAS 90.62% <ø> (ø) Carriedforward from df16ab4
hipBLASLt 68.77% <ø> (-1.30%) ⬇️ Carriedforward from df16ab4
hipCUB 82.68% <ø> (ø) Carriedforward from df16ab4
hipDNN 87.36% <100.00%> (+0.02%) ⬆️
hipFFT 42.47% <ø> (-0.18%) ⬇️ Carriedforward from df16ab4
hipRAND 76.12% <ø> (ø) Carriedforward from df16ab4
hipSOLVER 69.67% <ø> (ø) Carriedforward from df16ab4
hipSPARSE 87.03% <ø> (ø) Carriedforward from df16ab4
rocALUTION ?
rocBLAS 69.20% <ø> (ø) Carriedforward from df16ab4
rocFFT 49.54% <ø> (+3.05%) ⬆️ Carriedforward from df16ab4
rocRAND 57.32% <ø> (ø) Carriedforward from df16ab4
rocSOLVER 76.74% <ø> (+0.08%) ⬆️ Carriedforward from df16ab4
rocSPARSE 83.54% <ø> (-0.07%) ⬇️ Carriedforward from df16ab4
rocThrust 93.15% <ø> (ø) Carriedforward from df16ab4

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines Coverage Δ
...nn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp 58.61% <100.00%> (+1.47%) ⬆️

... and 273 files with indirect coverage changes

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
  • 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

…IOPEN-2684/miopen-provider-route-pooling-fwd

# Conflicts:
#	dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md
#	dnn-providers/miopen-provider/test_categories_integration.yaml
…IOPEN-2684/miopen-provider-route-pooling-fwd

# Conflicts:
#	dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md
adickin-amd added a commit that referenced this pull request Nov 6, 2026
… claim PRs (#15293)

## Summary
Stop regenerating `SUPPORT_MATRIX_OVERVIEW.md` in every PR that changes
bundles or support claims. Every such PR edited the same generated file,
so each merge made the other open claim PRs conflict on it and forced
another develop merge and regenerate (for example #14984 after #14259
and #14622 merged). The overview is now refreshed by hand in a PR of its
own, and no check fails when it lags.

JIRA ID : ALMIOPEN-3185

## Risk Assessment
Risk 1. Python tooling, a pre-commit hook and docs only; no product
code, no build change. The only behavior removed is the staleness check
on a generated file. Errors in the overview renderer and in the sidecars
it reads are still caught by the render hook and the
`hipdnn_support_features_python_tests` ctest.

## ASIC Coverage
None. Host-only Python tooling and docs; no GPU code path changes.

## Testing Summary
- Unit: `test_harvest_support_claims.py` and `test_support_features.py`,
including the real-tree render and size checks.
- Smoke: with the committed overview deliberately edited out of date,
the `TestRealBundleTree` ctest tests pass and `feature_matrix.py
--check` reports it as stale (exit 1).
- Pre-commit on the changed files, including the new
`render-support-matrix-overview` hook.

## Testing Checklist
- [x] Support-claim tooling tests - `python -m pytest
dnn-providers/integration-tests/scripts/test_harvest_support_claims.py
dnn-providers/integration-tests/scripts/test_support_features.py` -
Status: Passed (1914 tests)
- [x] Stale overview no longer fails - `pytest test_support_features.py
-k RealBundleTree` against an edited overview - Status: Passed
- [x] Pre-commit - `pre-commit run --files <changed files>` - Status:
Passed
- [ ] PR CI - GitHub PR checks - Status: Pending

## Technical Changes
- `test_support_features.py`: remove
`test_committed_overview_is_current`. The real-tree render and the size
limit stay.
- `.pre-commit-config.yaml`: replace `check-support-matrix-overview`
(`--check`) with `render-support-matrix-overview`, which renders to
stdout and discards, like `render-support-matrix`.
- `harvest_support_claims.py`: `--write` no longer re-renders the
overview; `refresh_overview` and its tests are removed, and a test
checks that `--write` leaves the overview alone.
- `feature_matrix.py`: the rendered overview states that it can lag
develop and points to the current copy in the `hipdnn-support-matrix` CI
artifact. `--check` stays for telling whether a refresh is due.
- `SUPPORT_MATRIX_OVERVIEW.md`: regenerated for that header note.
- `docs/support-matrix.md`, `docs/adding-tests.md`, RFC 0015: the
overview is not committed with sidecar or bundle changes; after a
harvest or a batch of claim PRs merges, whoever landed them refreshes it
in a PR with only that file.

Open claim PRs that already carry an overview change should drop it once
this merges: `git checkout origin/develop --
dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md`.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant