[AIROCMLIR-1078] Port MIOpen conv layout translation to rocmlirTriton

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)."""
 

mlir/utils/performance/perfRunner.py

@@ -134,6 +134,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
@@ -987,20 +1042,35 @@ def benchmark_external(cls, commandline, paths: Paths, arch, num_cu, num_chiplet
         config = cls.from_command_line(commandline, arch, num_cu, num_chiplets)
         config_args, _ = extract_tuning_key_metadata(commandline)
         config_args = drop_perf_priority(config_args)
-        miopen_driver_cmd = [MIOPENDRIVER, *config_args, '-V', '0', '-t', '1']
-        print("Running MIOpen Benchmark: ", ' '.join(config_args))
+        # 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(config_args)
+        if miopen_commandline is None:
+            print("Skipping MIOpen benchmark: conv layout has no equivalent MIOpen "
+                  f"NCHW/NHWC representation: {' '.join(config_args)}")
+            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)