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).""" diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py index 15d7dc87d4..c562d9260b 100755 --- a/mlir/utils/performance/perfRunner.py +++ b/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)