[AIROCMLIR-426] Translate rocMLIR conv layouts to MIOpen or skip#2422
Conversation
Signed-off-by: bogdan-petkovic <[email protected]>
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #2422 +/- ##
===========================================
+ Coverage 82.57% 82.70% +0.13%
===========================================
Files 120 120
Lines 42852 42828 -24
Branches 7110 7106 -4
===========================================
+ Hits 35381 35417 +36
+ Misses 4815 4794 -21
+ Partials 2656 2617 -39
Flags with carried forward coverage won't be shown. Click here to find out more. 🚀 New features to boost your workflow:
|
…2-miopen-layout-translate-or-skip Co-authored-by: Cursor <[email protected]> # Conflicts: # mlir/utils/performance/perfRunner.py
|
Ran this through nightly and compared against develop.
Both boards green |
There was a problem hiding this comment.
Pull request overview
This PR fixes MIOpen convolution benchmarking by translating rocMLIR-specific conv layout strings (e.g., GNC01, NGC01) into MIOpenDriver-supported NCHW/NHWC layouts when the mapping is exact, and skipping configs that cannot be represented fairly. It also tightens failure handling so MIOpen benchmark/parse failures raise errors instead of silently producing NaN, improving CI signal.
Changes:
- Add
rocmlir_layout_to_miopen()andconv_commandline_to_miopen_layouts()to translate-f/-I/-Olayouts toNCHW/NHWCor skip non-representable/mixed-layout configs. - Make MIOpen benchmark failures and elapsed-time parse failures raise
RuntimeErrorinstead of returningNaN. - Add unit tests for layout translation and translate-or-skip behavior.
Reviewed changes
Copilot reviewed 2 out of 2 changed files in this pull request and generated 2 comments.
| File | Description |
|---|---|
| mlir/utils/performance/perfRunner.py | Adds rocMLIR→MIOpen layout translation/skip logic and hard-fails on MIOpen benchmark/parse errors. |
| mlir/utils/performance/tests/test_perfRunner.py | Adds unit tests covering layout mapping and commandline translate-or-skip behavior. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
There was a problem hiding this comment.
Verdict: APPROVE -- submitted as COMMENT (automated reviews are advisory) · Findings: 0 (0 Critical, 0 Major, 0 Minor)
Scope
Python perf-tooling change in mlir/utils/performance/perfRunner.py (+ tests). Adds rocmlir_layout_to_miopen() / conv_commandline_to_miopen_layouts() to translate rocMLIR conv layout names (e.g. GNC01) into MIOpen's NCHW/NHWC — dropping the group dim G, renaming spatial 0->H/1->W, treating output channel K as C — or skipping configs with no faithful MIOpen equivalent (unrepresentable ordering, or filter/input/output disagreeing on NCHW vs NHWC). Also replaces silent-NaN MIOpen benchmark/parse failures with RuntimeError so CI fails on real errors.
Findings
No blocking issues found.
Notes
- Verified the translate-or-skip logic against the tensor letter conventions (
FILTER_LAYOUT_MAP/INPUT_LAYOUT_MAP/OUTPUT_LAYOUT_MAP): filter usesN+C, output usesN+K, input usesN+C, so theK->Creplacement never produces a double-Ccollision. The chainedstr.replaceorder is collision-free. - Confirmed the new error handling matches
run_pipeline's(outs, ok)contract (if not noerr:→ failure), and thatrun_pipelinealready prints the failing stage's stderr, so the "see the MIOpenDriver error above" wording in the raisedRuntimeErroris accurate. - Tests cover positive translation (
NCHW/NHWC), MIOpen pass-through,K-as-C, group conv (-gpreserved), unrepresentable ordering, and mixedNCHW/NHWCskip — good positive/negative coverage for a helper of this kind. - Minor observation (not blocking, no action needed):
conv_commandline_to_miopen_layoutsreturns the commandline unchanged if no-f/-I/-Oflags are present (emptyseen_layouts), but a conv config always carries these, so this defensive path is unreachable in practice.
CI status
No failing or cancelled checks. All functional checks are still pending/in-progress; detect passed. The auto-review pipeline's own review check being in-progress is expected and is not a CI failure.
justinrosner
left a comment
There was a problem hiding this comment.
When this gets all of the required approvals can you also create a PR porting this change over to rocmlirTriton?
There was a problem hiding this comment.
Why not make these helper functions members of ConvConfiguration similar to how benchmark_external is already setup?
There was a problem hiding this comment.
I put these right next to the other layout helpers on purpose. input_layouts, filter_layouts and the inverse ones are all plain module level functions too, and they're actually shared between ConvConfiguration and ConvGemmConfiguration, which is why they don't live inside a single class. benchmark_external is a classmethod because it genuinely needs cls and things like table_entry and from_command_line. These two new ones are just pure string helpers with no class state, so making them staticmethods would mostly add a class prefix while making them inconsistent with the helpers sitting right above them. I'm happy to move them into ConvConfiguration if you'd prefer, I just wanted to keep it consistent with what's already there
There was a problem hiding this comment.
Ah okay. I'm okay with leaving them next to the other layout helpers then.
| if len(seen_layouts) > 1: | ||
| return None |
There was a problem hiding this comment.
Is this really true ? can you double check ?
There was a problem hiding this comment.
I double-checked directly with MIOpenDriver on a gfx942 GPU, same conv, only changing the layouts:
# all NCHW
$ MIOpenDriver conv -F 1 -f NCHW -I NCHW -O NCHW -n 1 -c 512 -H 32 -W 32 -k 512 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -t 1 -V 0 -t 1
Elapsed: 0.113675 ms
# all NHWC
$ MIOpenDriver conv -F 1 -f NHWC -I NHWC -O NHWC ...
Elapsed: 0.109906 ms
# mixed: -f NCHW -I NCHW -O NHWC
$ MIOpenDriver conv -F 1 -f NCHW -I NCHW -O NHWC ...
No suitable algorithm was found to execute the required convolution (rc = 0x7)
Same config runs fine for all-NCHW and all-NHWC, but fails the moment the layouts are mixed, MIOpen has no solver for mixed filter/input/output layouts. That's why I skip those configs so if we translated and ran them, the fail-on-error path would turn the nightly red
umangyadav
left a comment
There was a problem hiding this comment.
Address justin's comment before merging
Motivation
While running the MIOpen benchmark, every conv config fails with
Invalid Layout Parameter Valuebecause rocMLIR passes its own layout names (e.g.GNC01,NGC01) toMIOpenDriver, which only acceptsNCHW/NHWC.The failures were silently converted to
NaN, so nightly CI has been green despite the MIOpen benchmarks never running.Resolves https://github.com/ROCm/rocMLIR-internal/issues/2212 (AIROCMLIR-426).
Supersedes #2298.
Technical Details
rocmlir_layout_to_miopen()andconv_commandline_to_miopen_layouts(): translate-f/-I/-OtoNCHW/NHWConly when the layout maps exactly (drop the group dimG— MIOpen conveys the group count via-g— and rename spatial dims0->H,1->W). Configs whose layouts have no faithful MIOpen equivalent, or where filter/input/output do not share a single layout, are skipped instead of benchmarked against a different layout.RuntimeErrorinstead of silently returningNaN, so CI fails on real errors.run_pipelinenow returns the failing process's stderr, so the error message is no longer empty.Over
tier1-conv-configs: 1296/1363 configs are translated and benchmarked, 67 are skipped (all genuinely non-representable orderings or mixed NCHW/NHWC).Test Plan
yapf/flake8clean on changed filespytest mlir/utils/performance/tests/test_perfRunner.pyTest Result
Submission Checklist