From 692d71d4664cc3eaf29190ccab1005a7f009e9f8 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <> Date: Mon, 9 Jun 2025 16:40:48 +0000 Subject: [PATCH] Add runner script for the Fremont shapes. --- .gitignore | 3 + script/convert_miopen_driver_to_profiler.py | 41 ++++--- script/run_conv_profiler.py | 113 ++++++++++++++++++++ 3 files changed, 144 insertions(+), 13 deletions(-) create mode 100644 script/run_conv_profiler.py diff --git a/.gitignore b/.gitignore index 599ef99e35..383f7c0d35 100644 --- a/.gitignore +++ b/.gitignore @@ -68,3 +68,6 @@ build*/ # Python cache __pycache__/ + +# Test data +test_data/* diff --git a/script/convert_miopen_driver_to_profiler.py b/script/convert_miopen_driver_to_profiler.py index 2ddcbb67cd..940dcce201 100644 --- a/script/convert_miopen_driver_to_profiler.py +++ b/script/convert_miopen_driver_to_profiler.py @@ -102,8 +102,7 @@ def add_conv_params_to_cmd(args, cmd): print('Not supported spatial dim (supported: 1, 2, 3)') exit(1) - -def run_ck_grouped_conv_fwd(args): +def get_ck_grouped_conv_fwd_cmd(args): args.ck_profier_op = "grouped_conv_fwd" parse_data_type(args) parse_layouts(args) @@ -119,15 +118,18 @@ def run_ck_grouped_conv_fwd(args): cmd += [str(args.in_channels)] add_conv_params_to_cmd(args, cmd) + return cmd + +def run_ck_grouped_conv_fwd(args): + cmd = get_ck_grouped_conv_fwd_cmd(args) run_ck_profiler_cmd(cmd) - -def run_ck_grouped_conv_bwd_data(args): +def get_ck_grouped_conv_bwd_data_cmd(args): args.ck_profier_op = "grouped_conv_bwd_data" parse_data_type(args) parse_layouts(args) - # Test all split K value from the list {1, 2, 4, 8, 32, 64, 128} - args.split_k_value = -1 + # Test all split K value from the list {-1, 1, 2, 4, 8, 32, 64, 128} + args.split_k_value = "all" cmd = [str(args.ck_profiler_cmd), str(args.ck_profier_op)] cmd += [str(args.data_type), str(args.layout)] @@ -139,15 +141,19 @@ def run_ck_grouped_conv_bwd_data(args): add_conv_params_to_cmd(args, cmd) cmd += [str(args.split_k_value)] + + return cmd + +def run_ck_grouped_conv_bwd_data(args): + cmd = get_ck_grouped_conv_bwd_data_cmd(args) run_ck_profiler_cmd(cmd) - -def run_ck_grouped_conv_bwd_weight(args): +def get_ck_grouped_conv_bwd_weight_cmd(args): args.ck_profier_op = "grouped_conv_bwd_weight" parse_data_type(args) parse_layouts(args) - # Test all split K value from the list {1, 2, 4, 8, 32, 64, 128} - args.split_k_value = -1 + # Test all split K value from the list {-1, 1, 2, 4, 8, 32, 64, 128} + args.split_k_value = "all" cmd = [str(args.ck_profiler_cmd), str(args.ck_profier_op)] cmd += [str(args.data_type), str(args.layout)] @@ -159,6 +165,11 @@ def run_ck_grouped_conv_bwd_weight(args): add_conv_params_to_cmd(args, cmd) cmd += [str(args.split_k_value)] + + return cmd + +def run_ck_grouped_conv_bwd_weight(args): + cmd = get_ck_grouped_conv_bwd_weight_cmd(args) run_ck_profiler_cmd(cmd) # Get name of miopen driver, remove it from unknown @@ -194,8 +205,7 @@ def run_ck_profiler(args): if args.forw == 0 or args.forw == 4 or args.forw == 5 or args.forw == 6: run_ck_grouped_conv_bwd_weight(args) - -if __name__ == "__main__": +def get_parser(): parser = argparse.ArgumentParser( prog="converter", description="Convert miopen driver command to ck Profiler" @@ -206,7 +216,7 @@ if __name__ == "__main__": "32 -F 1 -t 1", ) parser.add_argument( - "-in_layout", + "--in_layout", "-I", default="NCHW", type=str, @@ -405,6 +415,11 @@ if __name__ == "__main__": help="Number of Groups (Default=1)" ) + return parser + +if __name__ == "__main__": + parser = get_parser() + args, unknown = parser.parse_known_args() init_const_args(args) process_miopen_driver_name(args, unknown) diff --git a/script/run_conv_profiler.py b/script/run_conv_profiler.py new file mode 100644 index 0000000000..5b211ba360 --- /dev/null +++ b/script/run_conv_profiler.py @@ -0,0 +1,113 @@ +#!/usr/bin/env python3 + +import os +import argparse +import subprocess +import sys +import pandas as pd +from convert_miopen_driver_to_profiler import get_parser, init_const_args, process_miopen_driver_name, \ + get_ck_grouped_conv_fwd_cmd, get_ck_grouped_conv_bwd_data_cmd, get_ck_grouped_conv_bwd_weight_cmd + +def parse_cli_args(): + """Parse command line arguments""" + parser = argparse.ArgumentParser(description="Run CK convolution profiler.") + parser.add_argument("--csv-file", type=str, dest="csv_file", required=True, help="Path to the CSV file containing test cases.") + parser.add_argument("--fwd-only", action="store_true", help="Run only forward convolution.") + parser.add_argument("--bwd-data-only", action="store_true", help="Run only backward data convolution.") + parser.add_argument("--bwd-weight-only", action="store_true", help="Run only backward weight convolution.") + parser.add_argument("--start", type=int, default=None, help="Start index for processing shapes in the CSV file.") + parser.add_argument("--end", type=int, default=None, help="End index for processing shapes in the CSV file. If None, process all shapes.") + parser.add_argument("--no-verification", action="store_true", help="Disable verification in the CK profiler.") + + args, unknown_args = parser.parse_known_args() + + if unknown_args: + print(f"Unknown arguments: {unknown_args}", file=sys.stderr) + sys.exit(1) + + return args + +def parse_profiler_command(args, fwd_only=False, bwd_data_only=False, bwd_weight_only=False): + # MIOpen get number of channel per all groups, CK profiler get number of + # channel per group + args.in_channels = int(args.in_channels / args.group_count) + args.out_channels = int(args.out_channels / args.group_count) + + cmd = None + if fwd_only: + args.forw = 1 + cmd = get_ck_grouped_conv_fwd_cmd(args) + if bwd_data_only: + args.forw = 2 + cmd = get_ck_grouped_conv_bwd_data_cmd(args) + if bwd_weight_only: + args.forw = 4 + cmd = get_ck_grouped_conv_bwd_weight_cmd(args) + + return cmd + +def get_profiler_commands(csv_file, no_verification=False, fwd_only=False, bwd_data_only=False, bwd_weight_only=False): + if not os.path.isfile(csv_file): + print(f"Error: The specified CSV file '{csv_file}' does not exist.", file=sys.stderr) + sys.exit(1) + + df = pd.read_csv(csv_file) + shapes = df['Shape'].tolist() + parser = get_parser() + commands = [] + + for i, line in enumerate(shapes): + try: + args, unknown = parser.parse_known_args(line.split()) + + init_const_args(args) + process_miopen_driver_name(args, unknown) + assert len(unknown) == 4 and unknown[0] == "--fil_layout" and unknown[2] == "--out_layout" and unknown[1] == unknown[3], \ + f"Error: Unknown arguments do not match: {unknown}" + assert unknown[1] == args.in_layout, \ + f"Error: Input layout does not match unknown arguments: {unknown[1]} != {args.in_layout}" + + if no_verification: + args.verify = 0 + + # Ensure we run always the timing. + args.time = 1 + + command = parse_profiler_command(args, + fwd_only=fwd_only, + bwd_data_only=bwd_data_only, + bwd_weight_only=bwd_weight_only) + if command is not None: + + commands.append(command) + except AttributeError as e: + print(f"Error processing line {i}: {line}. Skipping the line.") + continue + + return commands + +def run_ck_profiler_cmd(cmd): + cmd_concatenated_str = "" + for arg in cmd: + cmd_concatenated_str += arg + " " + env_vars = os.environ.copy() + env_vars["CK_PROFILER_DISABLED_OPS"] = "DeviceGroupedConvBwdWeight_Dl;DeviceGroupedConvBwdWeight_Xdl_CShuffleV3" + env_vars["CK_PROFILER_OUTPUT_FILE"] = "/tmp/ck_profiler_output.txt" + subprocess.run(cmd,env=env_vars) + +def main(): + args = parse_cli_args() + profiler_commands = get_profiler_commands(args.csv_file, args.no_verification, args.fwd_only, args.bwd_data_only, args.bwd_weight_only) + print(f"Got {len(profiler_commands)} commands in total to run.") + if args.start is not None: + end = len(profiler_commands) + if args.end is not None: + end = min(args.end, end) + profiler_commands = profiler_commands[args.start-1:end] + + for i, cmd in enumerate(profiler_commands): + print(f"Running command {i + 1}/{len(profiler_commands)}: {cmd}") + run_ck_profiler_cmd(cmd) + +if __name__ == "__main__": + main() \ No newline at end of file