From 518d88faccf0ed292096a5260e540424588168cf Mon Sep 17 00:00:00 2001 From: Djordje Antic Date: Fri, 13 Mar 2026 15:55:37 -0400 Subject: [PATCH 1/3] [AIROCMLIR-426] rocMLIR - MIOpen layout translation --- mlir/utils/performance/perfRunner.py | 59 ++++++++++++++++++++++------ 1 file changed, 48 insertions(+), 11 deletions(-) diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py index 90b9010d27a5..6f1d8ec8cf19 100644 --- a/mlir/utils/performance/perfRunner.py +++ b/mlir/utils/performance/perfRunner.py @@ -95,6 +95,37 @@ def inverse_filter_layouts(filter_layout): return "".join(map[char] for char in filter_layout) +# Map rocMLIR-specific layout names to MIOpenDriver layout names (NCHW, NHWC). +# MIOpenDriver does not accept rocMLIR layout names (e.g. GNC01, NGC01). +ROCMLIR_TO_MIOPEN_LAYOUT = { + 'GNC01': 'NCHW', + 'NGC01': 'NCHW', + 'NC0G1': 'NCHW', + 'G0NC1': 'NCHW', + '01NGC': 'NHWC', + 'N01GC': 'NHWC', + 'N01GK': 'NHWC', + 'NCHW': 'NCHW', + 'NHWC': 'NHWC', +} + + +def conv_commandline_to_miopen_layouts(commandline): + """Return a copy of commandline with -f, -I, -O layout values translated to MIOpen names.""" + result = list(commandline) + for i in range(len(result)): + if result[i] == '-f' and i + 1 < len(result): + layout = result[i + 1] + result[i + 1] = ROCMLIR_TO_MIOPEN_LAYOUT.get(layout, layout) + elif result[i] == '-I' and i + 1 < len(result): + layout = result[i + 1] + result[i + 1] = ROCMLIR_TO_MIOPEN_LAYOUT.get(layout, layout) + elif result[i] == '-O' and i + 1 < len(result): + layout = result[i + 1] + result[i + 1] = ROCMLIR_TO_MIOPEN_LAYOUT.get(layout, layout) + return result + + @dataclass class MLIRPaths: rocmlir_gen_path: str @@ -694,20 +725,26 @@ 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'] + # Configs use rocMLIR layout names; MIOpenDriver expects NCHW/NHWC. + miopen_commandline = conv_commandline_to_miopen_layouts(commandline) + miopen_driver_cmd = [MIOPENDRIVER, *miopen_commandline, '-V', '0', '-t', '1'] print("Running MIOpen Benchmark: ", ' '.join(commandline)) # 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: + err_msg = outs.decode('utf-8') if isinstance(outs, bytes) else str(outs) + raise RuntimeError( + "MIOpen benchmark failed. CI must fail on MIOpen errors.\n" + "Failing command: " + ' '.join(miopen_driver_cmd) + "\n" + "Error: " + err_msg + ) + # 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 return config.table_entry(nanoseconds) From ea1b9d12145002eafcc37c3a6602e463992ea8d5 Mon Sep 17 00:00:00 2001 From: Cursor Agent Date: Sat, 14 Mar 2026 06:35:15 +0000 Subject: [PATCH 2/3] Fix yapf formatting in perfRunner.py Apply yapf-compliant formatting to raise RuntimeError() call to fix the py-checks CI failure. Co-authored-by: Djordje Antic --- mlir/utils/performance/perfRunner.py | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py index 6f1d8ec8cf19..731994a848e7 100644 --- a/mlir/utils/performance/perfRunner.py +++ b/mlir/utils/performance/perfRunner.py @@ -733,11 +733,9 @@ def benchmark_external(cls, commandline, paths: Paths, arch, num_cu, num_chiplet outs, noerr = run_pipeline([miopen_driver_cmd]) if not noerr: err_msg = outs.decode('utf-8') if isinstance(outs, bytes) else str(outs) - raise RuntimeError( - "MIOpen benchmark failed. CI must fail on MIOpen errors.\n" - "Failing command: " + ' '.join(miopen_driver_cmd) + "\n" - "Error: " + err_msg - ) + raise RuntimeError("MIOpen benchmark failed. CI must fail on MIOpen errors.\n" + "Failing command: " + ' '.join(miopen_driver_cmd) + "\n" + "Error: " + err_msg) # convert bytes to str outs = outs.decode('utf-8') # Extract Elapsed time in ms from the output of MIOpenDriver From d27c5d1abb32027f469eb3a34b2f56ff79fb6624 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=C4=90or=C4=91e=20Anti=C4=87?= Date: Sat, 14 Mar 2026 11:14:04 +0100 Subject: [PATCH 3/3] Collapsing nearly identical branches into a single check Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> --- mlir/utils/performance/perfRunner.py | 8 +------- 1 file changed, 1 insertion(+), 7 deletions(-) diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py index 731994a848e7..8f605e6a9ac8 100644 --- a/mlir/utils/performance/perfRunner.py +++ b/mlir/utils/performance/perfRunner.py @@ -114,13 +114,7 @@ def conv_commandline_to_miopen_layouts(commandline): """Return a copy of commandline with -f, -I, -O layout values translated to MIOpen names.""" result = list(commandline) for i in range(len(result)): - if result[i] == '-f' and i + 1 < len(result): - layout = result[i + 1] - result[i + 1] = ROCMLIR_TO_MIOPEN_LAYOUT.get(layout, layout) - elif result[i] == '-I' and i + 1 < len(result): - layout = result[i + 1] - result[i + 1] = ROCMLIR_TO_MIOPEN_LAYOUT.get(layout, layout) - elif result[i] == '-O' and i + 1 < len(result): + if result[i] in ('-f', '-I', '-O') and i + 1 < len(result): layout = result[i + 1] result[i + 1] = ROCMLIR_TO_MIOPEN_LAYOUT.get(layout, layout) return result