diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py index eb15ba4acd3f..2151111f1350 100644 --- a/mlir/utils/performance/perfRunner.py +++ b/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) diff --git a/mlir/utils/performance/tests/test_perfRunner.py b/mlir/utils/performance/tests/test_perfRunner.py index 6bb303b75974..66898c6dbe23 100644 --- a/mlir/utils/performance/tests/test_perfRunner.py +++ b/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)."""