Repository navigation
[AIROCMLIR-1078] Port MIOpen conv layout translation to rocmlirTriton - #359
Conversation
…il CI on MIOpen Signed-off-by: bogdan-petkovic <bpetkovi@amd.com>
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
Ports the MIOpen conv-layout translation fix from ROCm/rocMLIR PR #2422 to mlir/utils/performance/perfRunner.py. Adds rocmlir_layout_to_miopen() / conv_commandline_to_miopen_layouts() to translate -f/-I/-O layout args to NCHW/NHWC (or skip when non-representable / mixed), and makes ConvConfiguration.benchmark_external raise on MIOpen failure or output-parse failure instead of silently returning NaN.
Findings
No blocking issues found.
Notes
- The normalization in
rocmlir_layout_to_miopen(perfRunner.py:32) checks out against the on-command-line layout alphabet inFILTER_LAYOUT_MAP/INPUT_LAYOUT_MAP/OUTPUT_LAYOUT_MAP(perfRunner.py:85-87): e.g.GNC01→NCHW,NG01C→NHWC; thelen(seen_layouts) > 1guard (perfRunner.py:59) correctly skips inconsistent filter/input/output combinations. The.replace('K', 'C')step is effectively defensive (these configs express the output channel dim asC, notK) but is harmless. - The new
if not match:guard (perfRunner.py:879) fixes a latentAttributeErroronELAPSED_TIME_RE.search(...).group(1)when MIOpenDriver output can't be parsed — a real improvement over the prior code. - The skip path still returns
config.table_entry(np.nan)(perfRunner.py:863); this is an intentional "skip" signal (accompanied by a log line), distinct from the now-fatal genuine-error path, and matches the upstream rocMLIR behavior. - rocMLIR back-port:
mlir/utils/performance/is shared withROCm/rocMLIR, but the PR description explicitly identifies this as a direct port of rocMLIR PR #2422, satisfying the back-port note requirement. Unit tests were intentionally not ported (no perfRunner test harness in rocmlirTriton) — reasonable given the constraint.
CI status
No non-self checks are failing; py-checks/detect/ml-ci-internal.amd.com pass, Jenkins/Build and Test still pending. The auto-review check being in-progress is expected.
There was a problem hiding this comment.
Pull request overview
This PR updates the convolution performance runner to translate rocMLIR-style conv layout strings into the limited NCHW/NHWC layouts accepted by MIOpenDriver, skipping non-representable/mixed-layout configs and failing fast on real MIOpen benchmark or output-parse failures to avoid silently producing NaN.
Changes:
- Added
rocmlir_layout_to_miopen()andconv_commandline_to_miopen_layouts()to translate-f/-I/-Olayout arguments toNCHW/NHWCor skip when no faithful mapping exists. - Updated
ConvConfiguration.benchmark_external()to (a) skip configs without a faithful MIOpen layout and (b) raiseRuntimeErroron MIOpen execution failures or elapsed-time parse failures.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
You're right, we don't run the MIOpen benchmarking test at the moment. There is a ticket for it though, AIROCMLIR-901, which lists "Benchmark and Report Performance" as the remaining one to turn on. I picked it up from Djole last week. It depends on AIROCMLIR-899, because the performance checks need the tuning DB that the weekly job produces. The original 899 PR had gone stale and a lot of it already landed on develop, so I opened a fresh PR for just the remaining part #361. Once 899 is merged I'll do the 901 work to enable the nightly Benchmark and Report Performance stage, and then this port can go in on top of that. |
Would it be possible to add test harness ? |
rocmlirTriton does have perfRunner tests: see |
Signed-off-by: bogdan-petkovic <bpetkovi@amd.com>
added |
…anslate Signed-off-by: bogdan-petkovic <bpetkovi@amd.com>
…t-miopen-conv-layout-translate # Conflicts: # mlir/utils/performance/perfRunner.py
Signed-off-by: bogdan-petkovic <bpetkovi@amd.com>
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 toNaN, so nightly CI stayed green even though the MIOpen conv benchmarks never actually ran. This ports the fix from ROCm/rocMLIR#2422 over to rocmlirTriton.Technical Details
Adds
rocmlir_layout_to_miopen()andconv_commandline_to_miopen_layouts()tomlir/utils/performance/perfRunner.py, which translate the-f/-I/-Olayout arguments toNCHW/NHWConly when they map exactly (the group dimensionGis dropped since MIOpen conveys the group count via-g, and the spatial dims are renamed0->H,1->W). Configs whose layouts have no faithful MIOpen equivalent, or where filter/input/output do not agree on a singleNCHW/NHWClayout, are skipped instead of being benchmarked against a different layout.ConvConfiguration.benchmark_externalnow raisesRuntimeErroron a real MIOpen benchmark failure or an output-parse failure instead of silently returningNaN, so CI fails on genuine errors (the raised error points at MIOpenDriver's stderr, whichrun_pipelinealready prints to the log). This is a direct port of ROCm/rocMLIR#2422, unit tests were not ported because rocmlirTriton does not have a perfRunner test harness.Test Plan
yapfandflake8run clean on the changed file. The translate-or-skip logic was exercised locally over the fullmlir/utils/performance/configs/tier1-conv-configs, and the mixed-layout skip decision was confirmed directly againstMIOpenDriveron a gfx942 GPU (all-NCHW and all-NHWC run fine, a mixed filter/input/output layout fails at runtime with "No suitable algorithm was found to execute the required convolution").Test Result
tier1-conv-configs, 1296/1363 configs are translated and benchmarked (1294 NCHW, 2 NHWC) and 67 are skipped, all of which are genuinely non-representable orderings (G0NC1,01NGC) or mixedNCHW/NHWCcombinations. No layout errors and no silentNaNfallbacks remain.Submission Checklist