Skip to content

[AIROCMLIR-426] rocMLIR - MIOpen layout translation #2298

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
dorde-antic wants to merge 12 commits into
base: develop
Choose a base branch
from
+74 −12
Open

[AIROCMLIR-426] rocMLIR - MIOpen layout translation #2298

Changes from 2 commits
Commits
Show all changes
12 commits
Select commit Hold shift + click to select a range
518d88f
[AIROCMLIR-426] rocMLIR - MIOpen layout translation
dorde-antic Mar 13, 2026
fc3e55c
Merge branch 'develop' into miopen-benchmark-failures
dorde-antic Mar 13, 2026
ea1b9d1
Fix yapf formatting in perfRunner.py
cursoragent Mar 14, 2026
d27c5d1
Collapsing nearly identical branches into a single check
dorde-antic Mar 14, 2026
ea4fb03
Merge branch 'develop' into miopen-benchmark-failures
dorde-antic Mar 23, 2026
e8a6b67
Merge branch 'develop' into miopen-benchmark-failures
dorde-antic Mar 23, 2026
78b4bde
Apply suggestions from code review
dorde-antic Mar 24, 2026
58d5fba
Address comments
dorde-antic Apr 8, 2026
ceb8c23
Merge branch 'develop' into miopen-benchmark-failures
dorde-antic Apr 8, 2026
013436d
Address PR review comments
dorde-antic Apr 8, 2026
82fc60e
Merge branch 'develop' into miopen-benchmark-failures
dorde-antic Apr 15, 2026
a0631d0
Merge branch 'develop' into miopen-benchmark-failures
dorde-antic Apr 24, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
59 changes: 48 additions & 11 deletions mlir/utils/performance/perfRunner.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -95,6 +95,37 @@ def inverse_filter_layouts(filter_layout):
return "".join(map[char] for char in filter_layout)


# Map rocMLIR-specific layout names to MIOpenDriver layout names (NCHW, NHWC).
# MIOpenDriver does not accept rocMLIR layout names (e.g. GNC01, NGC01).

@umangyadav umangyadav Mar 24, 2026

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Add comment that this is "best effort mapping" and may not result in fair comparison.

Add comment that it is doing "channel first" mapping to NCHW and "channel last" to "NHWC"

ROCMLIR_TO_MIOPEN_LAYOUT = {

@dhernandez0 dhernandez0 Mar 26, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

we could simplify this by first converting 0 -> H and 1 -> W.

@dhernandez0 dhernandez0 Mar 26, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

also, I see we just drop "G", we can do that if it's not a group conv (G=1), if it is, we can't. Are we checking that somewhere?

@dhernandez0 dhernandez0 Mar 26, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

also input and output have "K" as well, we aren't converting that here? it's probably easier to keep separate dicts for input, filter and output.

@justinrosner justinrosner Apr 24, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should we emit an error (or at least a warning) if we are dropping G?

@umangyadav umangyadav Mar 26, 2026

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Did you run nightly reports to see if it runs into any errors or not ?

@dorde-antic dorde-antic Apr 8, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I've just kicked one https://ml-ci-internal.amd.com/job/MLIR/job/mlir/job/PR-2298/

'GNC01': 'NCHW',
'NGC01': 'NCHW',
'NC0G1': 'NCHW',
'G0NC1': 'NCHW',
'01NGC': 'NHWC',
'N01GC': 'NHWC',
'N01GK': 'NHWC',
Comment thread
dorde-antic marked this conversation as resolved.
Outdated
'NCHW': 'NCHW',
'NHWC': 'NHWC',
}

@cursor cursor Bot Mar 13, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It would be worth adding a comment (or an assert / logging.warning) for the fall-through case. Right now .get(layout, layout) silently passes through any unrecognised layout unchanged — which means a new rocMLIR layout name that's missing from this table will produce an 'Invalid Layout' error in MIOpenDriver that is hard to trace back here.

A lightweight safeguard:

def conv_commandline_to_miopen_layouts(commandline):
    ...
    if result[i] in layout_flags:
        layout = result[i + 1]
        translated = ROCMLIR_TO_MIOPEN_LAYOUT.get(layout)
        if translated is None:
            print(f"Warning: unknown rocMLIR layout '{layout}' passed through untranslated")
            translated = layout
        result[i + 1] = translated

Alternatively, document that ROCMLIR_TO_MIOPEN_LAYOUT is expected to be exhaustive and add a regression test covering any new layout names that arise.



def conv_commandline_to_miopen_layouts(commandline):
"""Return a copy of commandline with -f, -I, -O layout values translated to MIOpen names."""
result = list(commandline)
for i in range(len(result)):
if result[i] == '-f' and i + 1 < len(result):
layout = result[i + 1]
result[i + 1] = ROCMLIR_TO_MIOPEN_LAYOUT.get(layout, layout)
elif result[i] == '-I' and i + 1 < len(result):
layout = result[i + 1]
result[i + 1] = ROCMLIR_TO_MIOPEN_LAYOUT.get(layout, layout)
elif result[i] == '-O' and i + 1 < len(result):
layout = result[i + 1]
result[i + 1] = ROCMLIR_TO_MIOPEN_LAYOUT.get(layout, layout)

@cursor cursor Bot Mar 13, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The three if/elif branches are identical except for the flag name. You can collapse them into a single check with a set, which also makes it easy to add more layout flags in the future:

def conv_commandline_to_miopen_layouts(commandline):
    """Return a copy of commandline with -f, -I, -O layout values translated to MIOpen names."""
    result = list(commandline)
    layout_flags = {'-f', '-I', '-O'}
    for i in range(len(result) - 1):
        if result[i] in layout_flags:
            result[i + 1] = ROCMLIR_TO_MIOPEN_LAYOUT.get(result[i + 1], result[i + 1])
    return result

Also note that iterating range(len(result) - 1) removes the need for the repeated i + 1 < len(result) guard.

Comment thread
dorde-antic marked this conversation as resolved.
Outdated
return result
Comment on lines +98 to +148

Copilot AI Apr 24, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

New layout translation logic (_rocmlir_layout_to_miopen / conv_commandline_to_miopen_layouts) isn’t covered by the existing mlir/utils/performance/tests/test_perfRunner.py suite (which already tests layout helpers). Adding unit tests for representative layouts (e.g. GNC01, NGC01, grouped conv cases, and pass-through on unknown) would help prevent regressions and validate the heuristic.

Copilot uses AI. Check for mistakes.


@dataclass
class MLIRPaths:
rocmlir_gen_path: str
Expand Down Expand Up @@ -694,20 +725,26 @@ def benchmark_external(cls, commandline, paths: Paths, arch, num_cu, num_chiplet
if os.path.exists(get_profiler_output_path(arch, BENCHMARKING_METRICS_FILE_NAME)):
os.remove(get_profiler_output_path(arch, BENCHMARKING_METRICS_FILE_NAME))
config = cls.from_command_line(commandline, arch, num_cu, num_chiplets)
miopen_driver_cmd = [MIOPENDRIVER, *commandline, '-V', '0', '-t', '1']
# Configs use rocMLIR layout names; MIOpenDriver expects NCHW/NHWC.
miopen_commandline = conv_commandline_to_miopen_layouts(commandline)
miopen_driver_cmd = [MIOPENDRIVER, *miopen_commandline, '-V', '0', '-t', '1']
print("Running MIOpen Benchmark: ", ' '.join(commandline))

@cursor cursor Bot Mar 13, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The print still shows commandline (the original rocMLIR layout names) while the actual MIOpenDriver call uses miopen_commandline (with translated names). If a benchmark fails and someone tries to reproduce it by copy-pasting this line, they'll get the untranslated command and hit the same 'Invalid Layout' error all over again. Consider printing miopen_commandline instead, or at least the full miopen_driver_cmd:

print("Running MIOpen Benchmark: ", ' '.join(miopen_driver_cmd))

Comment thread
dorde-antic marked this conversation as resolved.
Outdated
# invoke MIOpenDriver.
outs, noerr = run_pipeline([miopen_driver_cmd])
nanoseconds = np.nan
if noerr:
# convert bytes to str
outs = outs.decode('utf-8')
# Extract Elapsed time in ms from the output of MIOpenDriver
# Use regular expression to match the contents between
# "Elasped: " (note the space at the end) and "ms"
elapsed_time_in_ms = ELAPSED_TIME_RE.search(outs).group(1)
nanoseconds = float(elapsed_time_in_ms) * 1.0e6

if not noerr:
err_msg = outs.decode('utf-8') if isinstance(outs, bytes) else str(outs)
raise RuntimeError(
"MIOpen benchmark failed. CI must fail on MIOpen errors.\n"
"Failing command: " + ' '.join(miopen_driver_cmd) + "\n"

@cursor cursor Bot Mar 13, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This raise RuntimeError(...) block is the cause of the py-checks CI failure. YAPF expects the first argument on the same line as raise RuntimeError( with subsequent lines aligned to match. The fix reformats this to:

raise RuntimeError("MIOpen benchmark failed. CI must fail on MIOpen errors.
"
                   "Failing command: " + ' '.join(miopen_driver_cmd) + "
"
                   "Error: " + err_msg)

Run yapf -i mlir/utils/performance/perfRunner.py locally to apply the fix.

"Error: " + err_msg

@cursor cursor Bot Mar 13, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

When noerr is False, outs holds the stdout of the failed process (from the p.communicate() call in run_pipeline's except block). The actual diagnostic output — the text that explains why MIOpenDriver failed — is almost always on stderr, which run_pipeline already captures and prints via p.stderr.read() before returning. So err_msg here will typically be empty or contain unrelated stdout.

Since run_pipeline already prints the failure details, you may simply want to raise without including err_msg, or restructure run_pipeline to surface stderr through its return value. At minimum, a note in the message that the real error was printed above would help:

if not noerr:
    raise RuntimeError(
        "MIOpen benchmark failed (see stderr output above for details).
"
        "Failing command: " + ' '.join(miopen_driver_cmd)
    )

)
# convert bytes to str
outs = outs.decode('utf-8')
# Extract Elapsed time in ms from the output of MIOpenDriver
# Use regular expression to match the contents between
# "Elasped: " (note the space at the end) and "ms"
Comment thread
dorde-antic marked this conversation as resolved.
Outdated
elapsed_time_in_ms = ELAPSED_TIME_RE.search(outs).group(1)
Comment thread
dorde-antic marked this conversation as resolved.
Outdated
nanoseconds = float(elapsed_time_in_ms) * 1.0e6
return config.table_entry(nanoseconds)


Expand Down
Loading