[AIROCMLIR-426] Translate rocMLIR conv layouts to MIOpen or skip

mlir/utils/performance/perfRunner.py

@@ -98,6 +98,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
@@ -828,20 +883,35 @@ def benchmark_external(cls, commandline, paths: Paths, arch, num_cu, num_chiplet
         if config.datatype not in cls.MIOPEN_SUPPORTED_DTYPES:
             print(f"Skipping MIOpen benchmark for unsupported datatype: {config.datatype}")
             return config.table_entry(np.nan)
-        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)
 
 

mlir/utils/performance/tests/test_perfRunner.py

@@ -217,6 +217,72 @@ def test_inverse_roundtrip(self):
         assert perfRunner.inverse_filter_layouts(perfRunner.filter_layouts(layout)) == layout
 
 
+class TestRocmlirLayoutToMiopen:
+    """Tests for rocmlir_layout_to_miopen (rocMLIR -> MIOpen layout mapping)."""
+
+    def test_channel_first_maps_to_nchw(self):
+        # G dropped, 0->H, 1->W; channel stays second.
+        assert perfRunner.rocmlir_layout_to_miopen("NGC01") == "NCHW"
+        assert perfRunner.rocmlir_layout_to_miopen("GNC01") == "NCHW"
+        assert perfRunner.rocmlir_layout_to_miopen("NC0G1") == "NCHW"
+
+    def test_channel_last_maps_to_nhwc(self):
+        assert perfRunner.rocmlir_layout_to_miopen("N01GC") == "NHWC"
+        assert perfRunner.rocmlir_layout_to_miopen("GN01C") == "NHWC"
+
+    def test_already_miopen_layouts_pass_through(self):
+        assert perfRunner.rocmlir_layout_to_miopen("NCHW") == "NCHW"
+        assert perfRunner.rocmlir_layout_to_miopen("NHWC") == "NHWC"
+
+    def test_output_channel_letter_k_treated_as_c(self):
+        assert perfRunner.rocmlir_layout_to_miopen("NGK01") == "NCHW"
+        assert perfRunner.rocmlir_layout_to_miopen("N01GK") == "NHWC"
+
+    def test_unrepresentable_orderings_return_none(self):
+        # These orderings (H/W or channel not first/last) have no NCHW/NHWC equivalent.
+        assert perfRunner.rocmlir_layout_to_miopen("G0NC1") is None
+        assert perfRunner.rocmlir_layout_to_miopen("01NGC") is None
+
+
+class TestConvCommandlineToMiopenLayouts:
+    """Tests for conv_commandline_to_miopen_layouts (translate-or-skip)."""
+
+    @staticmethod
+    def _cmd(f, i, o, group=1):
+        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=f, i=i, o=o,
+                                                                         g=group).split()
+
+    def test_consistent_nchw_config_is_translated(self):
+        result = perfRunner.conv_commandline_to_miopen_layouts(self._cmd("GNC01", "NGC01", "NGC01"))
+        assert result is not None
+        # Every -f/-I/-O value must now be NCHW.
+        for flag in ("-f", "-I", "-O"):
+            assert result[result.index(flag) + 1] == "NCHW"
+
+    def test_consistent_nhwc_config_is_translated(self):
+        result = perfRunner.conv_commandline_to_miopen_layouts(self._cmd("GN01C", "N01GC", "N01GC"))
+        assert result is not None
+        for flag in ("-f", "-I", "-O"):
+            assert result[result.index(flag) + 1] == "NHWC"
+
+    def test_group_conv_layout_is_still_translated(self):
+        # Dropping G from the layout string is valid; the group count rides on -g.
+        result = perfRunner.conv_commandline_to_miopen_layouts(
+            self._cmd("GNC01", "NGC01", "NGC01", group=2))
+        assert result is not None
+        assert result[result.index("-g") + 1] == "2"
+
+    def test_unrepresentable_layout_is_skipped(self):
+        assert perfRunner.conv_commandline_to_miopen_layouts(
+            self._cmd("G0NC1", "G0NC1", "NGC01", group=3)) is None
+
+    def test_mixed_nchw_nhwc_config_is_skipped(self):
+        # filter -> NCHW but output -> NHWC: no single MIOpen layout, so skip.
+        assert perfRunner.conv_commandline_to_miopen_layouts(self._cmd("GNC01", "NGC01",
+                                                                       "N01GC")) is None
+
+
 class TestGetNanoseconds:
     """Tests for get_nanoseconds (reads CSV from rocprof)."""