From 6a5fde37ce3e2fe2d5cc41ab8d0f825d531fefe1 Mon Sep 17 00:00:00 2001 From: bogdan-petkovic Date: Mon, 20 Jul 2026 07:15:16 -0500 Subject: [PATCH 1/3] [AIROCMLIR-1078] Translate rocMLIR conv layouts to MIOpen or skip, fail CI on MIOpen Signed-off-by: bogdan-petkovic --- mlir/utils/performance/perfRunner.py | 94 ++++++++++++++++++++++++---- 1 file changed, 82 insertions(+), 12 deletions(-) diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py index 894a38296b..bc65b0a20a 100755 --- a/mlir/utils/performance/perfRunner.py +++ b/mlir/utils/performance/perfRunner.py @@ -121,6 +121,61 @@ def inverse_filter_layouts(filter_layout): return "".join(map[char] for char in filter_layout) +# MIOpenDriver only understands the NCHW / NHWC memory layouts, whereas rocMLIR +# configs use richer names such as GNC01 / NGC01 that additionally encode the group +# dimension (G) and use 0/1 for the spatial dims. A config can therefore only be +# benchmarked against MIOpen when its layouts map *exactly* onto NCHW/NHWC once the +# group dimension is dropped (MIOpen conveys the group count separately via -g) and +# the spatial dims are renamed 0->H, 1->W. Layouts with any other ordering have no +# faithful MIOpen equivalent and are skipped instead of being benchmarked against a +# different layout (which would be an unfair comparison). +MIOPEN_CONV_LAYOUTS = {'NCHW', 'NHWC'} + + +def rocmlir_layout_to_miopen(layout): + """Map a rocMLIR conv layout name onto a MIOpenDriver layout, or None. + + The group dimension ``G`` is dropped (MIOpen passes the group count through the + separate ``-g`` flag) and the spatial dims are renamed (``0`` -> ``H``, ``1`` -> + ``W``). MIOpen spells every tensor's layout generically as NCHW/NHWC, so the + output tensor's channel letter ``K`` is treated like ``C``. + + Returns ``"NCHW"`` or ``"NHWC"`` when the layout is exactly one of those orderings, + otherwise ``None`` -- meaning the config is not MIOpen-representable and should be + skipped to keep the comparison fair. + """ + normalized = layout.replace('0', 'H').replace('1', 'W').replace('G', '').replace('K', 'C') + if normalized in MIOPEN_CONV_LAYOUTS: + return normalized + return None + + +def conv_commandline_to_miopen_layouts(commandline): + """Translate rocMLIR conv layout args (-f/-I/-O) into MIOpen layout names. + + Returns a new commandline list with the layout values replaced by their MIOpen + equivalents, or ``None`` when the configuration has no faithful MIOpen + representation -- either because a layout uses an ordering MIOpen cannot express, + or because the filter/input/output tensors do not share a single NCHW/NHWC layout. + Callers should skip the MIOpen benchmark in the ``None`` case rather than run an + unfair comparison. + """ + result = list(commandline) + layout_flags = {'-f', '-I', '-O'} + seen_layouts = set() + for i in range(len(result) - 1): + if result[i] in layout_flags: + miopen_layout = rocmlir_layout_to_miopen(result[i + 1]) + if miopen_layout is None: + return None + result[i + 1] = miopen_layout + seen_layouts.add(miopen_layout) + # MIOpen expects a single, consistent layout across filter, input and output. + if len(seen_layouts) > 1: + return None + return result + + @dataclass class MLIRPaths: rocmlir_gen_path: str @@ -798,20 +853,35 @@ 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'] - print("Running MIOpen Benchmark: ", ' '.join(commandline)) + # rocMLIR configs use layout names (e.g. GNC01) that MIOpenDriver rejects. + # Translate them to NCHW/NHWC; skip configs that have no faithful MIOpen + # equivalent instead of forcing an unfair comparison. + miopen_commandline = conv_commandline_to_miopen_layouts(commandline) + if miopen_commandline is None: + print("Skipping MIOpen benchmark: conv layout has no equivalent MIOpen " + f"NCHW/NHWC representation: {' '.join(commandline)}") + return config.table_entry(np.nan) + miopen_driver_cmd = [MIOPENDRIVER, *miopen_commandline, '-V', '0', '-t', '1'] + print("Running MIOpen Benchmark: ", ' '.join(miopen_driver_cmd)) # 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: + # run_pipeline already prints MIOpenDriver's stderr. A genuine MIOpen failure + # must fail CI instead of silently yielding NaN. + raise RuntimeError("MIOpen benchmark failed (see the MIOpenDriver error above); " + "CI must fail on MIOpen errors.\n" + f"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. Match the text + # between "Elapsed: " (note the trailing space) and "ms". + match = ELAPSED_TIME_RE.search(outs) + if not match: + raise RuntimeError("Failed to parse elapsed time from MIOpenDriver output.\n" + f"Failing command: {' '.join(miopen_driver_cmd)}\n" + f"Output:\n{outs}") + elapsed_time_in_ms = match.group(1) + nanoseconds = float(elapsed_time_in_ms) * 1.0e6 return config.table_entry(nanoseconds) From 0c66b1235e4a44624b6a0bc7ca8cca229b6888a4 Mon Sep 17 00:00:00 2001 From: bogdan-petkovic Date: Wed, 22 Jul 2026 11:09:51 +0000 Subject: [PATCH 2/3] Add perf-scripts lit test for conv layout translation Signed-off-by: bogdan-petkovic --- .../perf-scripts/perfRunner-conv-layout.py | 104 ++++++++++++++++++ 1 file changed, 104 insertions(+) create mode 100644 mlir/test/perf-scripts/perfRunner-conv-layout.py diff --git a/mlir/test/perf-scripts/perfRunner-conv-layout.py b/mlir/test/perf-scripts/perfRunner-conv-layout.py new file mode 100644 index 0000000000..fe54771207 --- /dev/null +++ b/mlir/test/perf-scripts/perfRunner-conv-layout.py @@ -0,0 +1,104 @@ +#!/usr/bin/env python3 +# +# Part of the MLIR Project, under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +"""Pure-Python coverage for the rocMLIR -> MIOpen conv layout translation. + +MIOpenDriver only understands NCHW/NHWC, while rocMLIR configs use richer layout +names (e.g. GNC01, NGC01). perfRunner translates the -f/-I/-O layouts to NCHW/NHWC +when they map exactly and skips configs with no faithful MIOpen equivalent. These +tests pin that translate-or-skip behaviour so it can't regress; they mirror the +unit tests added alongside the same change in rocMLIR (PR #2422). + +# RUN: %python %s +""" + +from pathlib import Path +import sys +import unittest + +MLIR_DIR = Path(__file__).resolve().parents[2] +PERF_DIR = MLIR_DIR / "utils" / "performance" +sys.path.insert(0, str(PERF_DIR)) + +from perfRunner import rocmlir_layout_to_miopen, conv_commandline_to_miopen_layouts # noqa: E402 + + +def make_conv_commandline(fil, inp, out, group=1): + """Build a minimal conv commandline (as a token list) with the given layouts.""" + return ("conv -F 1 -f {f} -I {i} -O {o} -n 1 -c 8 -H 16 -W 16 -k 8 " + "-y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -g {g}").format(f=fil, i=inp, o=out, + g=group).split() + + +class RocmlirLayoutToMiopenTest(unittest.TestCase): + """Tests for rocmlir_layout_to_miopen (single layout string -> MIOpen name).""" + + def test_channel_first_maps_to_nchw(self): + """Dropping G and renaming 0/1 leaves the channel second, i.e. NCHW.""" + self.assertEqual(rocmlir_layout_to_miopen("NGC01"), "NCHW") + self.assertEqual(rocmlir_layout_to_miopen("GNC01"), "NCHW") + self.assertEqual(rocmlir_layout_to_miopen("NC0G1"), "NCHW") + + def test_channel_last_maps_to_nhwc(self): + """A trailing channel dim maps to NHWC.""" + self.assertEqual(rocmlir_layout_to_miopen("N01GC"), "NHWC") + self.assertEqual(rocmlir_layout_to_miopen("GN01C"), "NHWC") + + def test_already_miopen_layouts_pass_through(self): + """NCHW/NHWC are returned unchanged.""" + self.assertEqual(rocmlir_layout_to_miopen("NCHW"), "NCHW") + self.assertEqual(rocmlir_layout_to_miopen("NHWC"), "NHWC") + + def test_output_channel_letter_k_treated_as_c(self): + """The output tensor spells the channel dim as K; MIOpen still wants NCHW/NHWC.""" + self.assertEqual(rocmlir_layout_to_miopen("NGK01"), "NCHW") + self.assertEqual(rocmlir_layout_to_miopen("N01GK"), "NHWC") + + def test_unrepresentable_orderings_return_none(self): + """Orderings that aren't NCHW/NHWC (channel or spatial in the wrong slot) skip.""" + self.assertIsNone(rocmlir_layout_to_miopen("G0NC1")) + self.assertIsNone(rocmlir_layout_to_miopen("01NGC")) + + +class ConvCommandlineToMiopenLayoutsTest(unittest.TestCase): + """Tests for conv_commandline_to_miopen_layouts (whole commandline translate-or-skip).""" + + def test_consistent_nchw_config_is_translated(self): + """A config whose filter/input/output all map to NCHW is translated.""" + result = conv_commandline_to_miopen_layouts(make_conv_commandline( + "GNC01", "NGC01", "NGC01")) + self.assertIsNotNone(result) + for flag in ("-f", "-I", "-O"): + self.assertEqual(result[result.index(flag) + 1], "NCHW") + + def test_consistent_nhwc_config_is_translated(self): + """A config whose filter/input/output all map to NHWC is translated.""" + result = conv_commandline_to_miopen_layouts(make_conv_commandline( + "GN01C", "N01GC", "N01GC")) + self.assertIsNotNone(result) + for flag in ("-f", "-I", "-O"): + self.assertEqual(result[result.index(flag) + 1], "NHWC") + + def test_group_conv_layout_is_still_translated(self): + """Dropping G from the layout is valid; the group count rides on -g.""" + result = conv_commandline_to_miopen_layouts( + make_conv_commandline("GNC01", "NGC01", "NGC01", group=2)) + self.assertIsNotNone(result) + self.assertEqual(result[result.index("-g") + 1], "2") + + def test_unrepresentable_layout_is_skipped(self): + """A layout with no NCHW/NHWC equivalent makes the whole config skip.""" + self.assertIsNone( + conv_commandline_to_miopen_layouts( + make_conv_commandline("G0NC1", "G0NC1", "NGC01", group=3))) + + def test_mixed_nchw_nhwc_config_is_skipped(self): + """MIOpen has no solver for mixed filter/input/output layouts, so skip.""" + self.assertIsNone( + conv_commandline_to_miopen_layouts(make_conv_commandline("GNC01", "NGC01", "N01GC"))) + + +if __name__ == "__main__": + unittest.main(argv=[sys.argv[0]], verbosity=2) From 234e365d30df337a99a8976a5a7e345da71206e6 Mon Sep 17 00:00:00 2001 From: bogdan-petkovic Date: Thu, 10 Sep 2026 08:03:46 +0000 Subject: [PATCH 3/3] Move conv layout tests into perfRunner-test.py Signed-off-by: bogdan-petkovic --- .../perf-scripts/perfRunner-conv-layout.py | 110 ------------------ mlir/test/perf-scripts/perfRunner-test.py | 81 +++++++++++++ 2 files changed, 81 insertions(+), 110 deletions(-) delete mode 100644 mlir/test/perf-scripts/perfRunner-conv-layout.py diff --git a/mlir/test/perf-scripts/perfRunner-conv-layout.py b/mlir/test/perf-scripts/perfRunner-conv-layout.py deleted file mode 100644 index bb2d1792a7..0000000000 --- a/mlir/test/perf-scripts/perfRunner-conv-layout.py +++ /dev/null @@ -1,110 +0,0 @@ -#!/usr/bin/env python3 -# -# Part of the MLIR Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -"""Pure-Python coverage for the rocMLIR -> MIOpen conv layout translation. - -MIOpenDriver only understands NCHW/NHWC, while rocMLIR configs use richer layout -names (e.g. GNC01, NGC01). perfRunner translates the -f/-I/-O layouts to NCHW/NHWC -when they map exactly and skips configs with no faithful MIOpen equivalent. These -tests pin that translate-or-skip behaviour so it can't regress; they mirror the -unit tests added alongside the same change in rocMLIR (PR #2422). - -# RUN: %python %s -""" - -import os -import shutil -import sys -import unittest - -# perfRunner.py depends on the compiled amd_arch_db binding. Both are deployed -# together by ci-performance-scripts, so import that copy rather than the source -# file, where the binding is unavailable. -_script = shutil.which("perfRunner.py") -if _script is None: - sys.exit("perfRunner.py not on PATH; did you run " - "`ninja ci-performance-scripts`?") -sys.path.insert(0, os.path.dirname(_script)) - -from perfRunner import rocmlir_layout_to_miopen, conv_commandline_to_miopen_layouts # noqa: E402 - - -def make_conv_commandline(fil, inp, out, group=1): - """Build a minimal conv commandline (as a token list) with the given layouts.""" - return ("conv -F 1 -f {f} -I {i} -O {o} -n 1 -c 8 -H 16 -W 16 -k 8 " - "-y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -g {g}").format(f=fil, i=inp, o=out, - g=group).split() - - -class RocmlirLayoutToMiopenTest(unittest.TestCase): - """Tests for rocmlir_layout_to_miopen (single layout string -> MIOpen name).""" - - def test_channel_first_maps_to_nchw(self): - """Dropping G and renaming 0/1 leaves the channel second, i.e. NCHW.""" - self.assertEqual(rocmlir_layout_to_miopen("NGC01"), "NCHW") - self.assertEqual(rocmlir_layout_to_miopen("GNC01"), "NCHW") - self.assertEqual(rocmlir_layout_to_miopen("NC0G1"), "NCHW") - - def test_channel_last_maps_to_nhwc(self): - """A trailing channel dim maps to NHWC.""" - self.assertEqual(rocmlir_layout_to_miopen("N01GC"), "NHWC") - self.assertEqual(rocmlir_layout_to_miopen("GN01C"), "NHWC") - - def test_already_miopen_layouts_pass_through(self): - """NCHW/NHWC are returned unchanged.""" - self.assertEqual(rocmlir_layout_to_miopen("NCHW"), "NCHW") - self.assertEqual(rocmlir_layout_to_miopen("NHWC"), "NHWC") - - def test_output_channel_letter_k_treated_as_c(self): - """The output tensor spells the channel dim as K; MIOpen still wants NCHW/NHWC.""" - self.assertEqual(rocmlir_layout_to_miopen("NGK01"), "NCHW") - self.assertEqual(rocmlir_layout_to_miopen("N01GK"), "NHWC") - - def test_unrepresentable_orderings_return_none(self): - """Orderings that aren't NCHW/NHWC (channel or spatial in the wrong slot) skip.""" - self.assertIsNone(rocmlir_layout_to_miopen("G0NC1")) - self.assertIsNone(rocmlir_layout_to_miopen("01NGC")) - - -class ConvCommandlineToMiopenLayoutsTest(unittest.TestCase): - """Tests for conv_commandline_to_miopen_layouts (whole commandline translate-or-skip).""" - - def test_consistent_nchw_config_is_translated(self): - """A config whose filter/input/output all map to NCHW is translated.""" - result = conv_commandline_to_miopen_layouts(make_conv_commandline( - "GNC01", "NGC01", "NGC01")) - self.assertIsNotNone(result) - for flag in ("-f", "-I", "-O"): - self.assertEqual(result[result.index(flag) + 1], "NCHW") - - def test_consistent_nhwc_config_is_translated(self): - """A config whose filter/input/output all map to NHWC is translated.""" - result = conv_commandline_to_miopen_layouts(make_conv_commandline( - "GN01C", "N01GC", "N01GC")) - self.assertIsNotNone(result) - for flag in ("-f", "-I", "-O"): - self.assertEqual(result[result.index(flag) + 1], "NHWC") - - def test_group_conv_layout_is_still_translated(self): - """Dropping G from the layout is valid; the group count rides on -g.""" - result = conv_commandline_to_miopen_layouts( - make_conv_commandline("GNC01", "NGC01", "NGC01", group=2)) - self.assertIsNotNone(result) - self.assertEqual(result[result.index("-g") + 1], "2") - - def test_unrepresentable_layout_is_skipped(self): - """A layout with no NCHW/NHWC equivalent makes the whole config skip.""" - self.assertIsNone( - conv_commandline_to_miopen_layouts( - make_conv_commandline("G0NC1", "G0NC1", "NGC01", group=3))) - - def test_mixed_nchw_nhwc_config_is_skipped(self): - """MIOpen has no solver for mixed filter/input/output layouts, so skip.""" - self.assertIsNone( - conv_commandline_to_miopen_layouts(make_conv_commandline("GNC01", "NGC01", "N01GC"))) - - -if __name__ == "__main__": - unittest.main(argv=[sys.argv[0]], verbosity=2) diff --git a/mlir/test/perf-scripts/perfRunner-test.py b/mlir/test/perf-scripts/perfRunner-test.py index 7ee947af1d..136ffae703 100644 --- a/mlir/test/perf-scripts/perfRunner-test.py +++ b/mlir/test/perf-scripts/perfRunner-test.py @@ -197,6 +197,87 @@ def test_inverse_roundtrip(self): layout) +def make_conv_commandline(fil, inp, out, group=1): + """Build a minimal conv commandline (as a token list) with the given layouts.""" + return ("conv -F 1 -f {f} -I {i} -O {o} -n 1 -c 8 -H 16 -W 16 -k 8 " + "-y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -g {g}").format(f=fil, i=inp, o=out, + g=group).split() + + +class RocmlirLayoutToMiopenTest(unittest.TestCase): + """Tests for rocmlir_layout_to_miopen (single layout string -> MIOpen name). + + MIOpenDriver only accepts NCHW/NHWC, so a rocMLIR layout is only usable once the + group dim is dropped (MIOpen passes the group count via -g) and the spatial dims + are renamed 0->H, 1->W. Anything else has no faithful MIOpen equivalent. + """ + + def test_channel_first_maps_to_nchw(self): + """Dropping G and renaming 0/1 leaves the channel second, i.e. NCHW.""" + self.assertEqual(perfRunner.rocmlir_layout_to_miopen("NGC01"), "NCHW") + self.assertEqual(perfRunner.rocmlir_layout_to_miopen("GNC01"), "NCHW") + self.assertEqual(perfRunner.rocmlir_layout_to_miopen("NC0G1"), "NCHW") + + def test_channel_last_maps_to_nhwc(self): + """A trailing channel dim maps to NHWC.""" + self.assertEqual(perfRunner.rocmlir_layout_to_miopen("N01GC"), "NHWC") + self.assertEqual(perfRunner.rocmlir_layout_to_miopen("GN01C"), "NHWC") + + def test_already_miopen_layouts_pass_through(self): + """NCHW/NHWC are returned unchanged.""" + self.assertEqual(perfRunner.rocmlir_layout_to_miopen("NCHW"), "NCHW") + self.assertEqual(perfRunner.rocmlir_layout_to_miopen("NHWC"), "NHWC") + + def test_output_channel_letter_k_treated_as_c(self): + """The output tensor spells the channel dim as K; MIOpen still wants NCHW/NHWC.""" + self.assertEqual(perfRunner.rocmlir_layout_to_miopen("NGK01"), "NCHW") + self.assertEqual(perfRunner.rocmlir_layout_to_miopen("N01GK"), "NHWC") + + def test_unrepresentable_orderings_return_none(self): + """Orderings that aren't NCHW/NHWC (channel or spatial in the wrong slot) skip.""" + self.assertIsNone(perfRunner.rocmlir_layout_to_miopen("G0NC1")) + self.assertIsNone(perfRunner.rocmlir_layout_to_miopen("01NGC")) + + +class ConvCommandlineToMiopenLayoutsTest(unittest.TestCase): + """Tests for conv_commandline_to_miopen_layouts (whole commandline translate-or-skip).""" + + def test_consistent_nchw_config_is_translated(self): + """A config whose filter/input/output all map to NCHW is translated.""" + result = perfRunner.conv_commandline_to_miopen_layouts( + make_conv_commandline("GNC01", "NGC01", "NGC01")) + self.assertIsNotNone(result) + for flag in ("-f", "-I", "-O"): + self.assertEqual(result[result.index(flag) + 1], "NCHW") + + def test_consistent_nhwc_config_is_translated(self): + """A config whose filter/input/output all map to NHWC is translated.""" + result = perfRunner.conv_commandline_to_miopen_layouts( + make_conv_commandline("GN01C", "N01GC", "N01GC")) + self.assertIsNotNone(result) + for flag in ("-f", "-I", "-O"): + self.assertEqual(result[result.index(flag) + 1], "NHWC") + + def test_group_conv_layout_is_still_translated(self): + """Dropping G from the layout is valid; the group count rides on -g.""" + result = perfRunner.conv_commandline_to_miopen_layouts( + make_conv_commandline("GNC01", "NGC01", "NGC01", group=2)) + self.assertIsNotNone(result) + self.assertEqual(result[result.index("-g") + 1], "2") + + def test_unrepresentable_layout_is_skipped(self): + """A layout with no NCHW/NHWC equivalent makes the whole config skip.""" + self.assertIsNone( + perfRunner.conv_commandline_to_miopen_layouts( + make_conv_commandline("G0NC1", "G0NC1", "NGC01", group=3))) + + def test_mixed_nchw_nhwc_config_is_skipped(self): + """MIOpen has no solver for mixed filter/input/output layouts, so skip.""" + self.assertIsNone( + perfRunner.conv_commandline_to_miopen_layouts( + make_conv_commandline("GNC01", "NGC01", "N01GC"))) + + class GetNanosecondsTest(TempFileTestCase): """Tests for get_nanoseconds (reads the CSV rocprof leaves behind)."""