[rocRoller] Add unique IDs to rrsuite

shared/rocroller/client/src/gemm.cpp

@@ -1071,14 +1071,14 @@ int main(int argc, const char* argv[])
     rocRoller::Client::GEMMClient::SolutionParameters solution{
         .macM = 64,
         .macN = 64,
-        .macK = 64,
+        .macK = -1,
 
         .waveM = -1,
         .waveN = -1,
         .waveK = -1,
         .waveB = -1,
 
-        .workgroupSizeX         = 128,
+        .workgroupSizeX         = -1,
         .workgroupSizeY         = 2,
         .workgroupMappingDim    = -1,
         .workgroupRemapXCC      = false,
@@ -1268,7 +1268,7 @@ int main(int argc, const char* argv[])
     app.add_option(
         "--mi",
         [&solution](auto& args) -> bool { return ParseMI(args[0], solution); },
-        "MI instruction to use.  Default 32x32x2x1 for floats, 32x32x8x1 for halfs.");
+        "MI instruction to use");
 
     app.add_option(
         "--workgroup_size_x", solution.workgroupSizeX, "Workgroup size in the x dimension.");
@@ -1527,6 +1527,11 @@ int main(int argc, const char* argv[])
         types.scaleBlockSize = arch.GetCapability(GPUCapability::DefaultScaleBlockSize);
     }
 
+    if(solution.workgroupSizeX == -1)
+        solution.workgroupSizeX = 2 * arch.GetCapability(GPUCapability::DefaultWavefrontSize);
+    if(solution.workgroupSizeY == -1)
+        solution.workgroupSizeY = 2;
+
     AssertFatal((types.typeAcc == "float") || (types.typeAcc == "half")
                 || (types.typeAcc == "bf16"));
 
@@ -1538,9 +1543,12 @@ int main(int argc, const char* argv[])
     io.doSaveAsm = asmOption->count() > 0;
     io.doSaveCO  = coOption->count() > 0;
 
-    // Set default MI sizes
+    // Set default MI and macK sizes
     if(arch.HasCapability(GPUCapability::HasMFMA))
     {
+        if(solution.macK == -1)
+            solution.macK = 64;
+
         if(types.typeA == "float" && types.typeB == "float" && types.typeC == "float"
            && types.typeD == "float")
         {
@@ -1592,6 +1600,9 @@ int main(int argc, const char* argv[])
     {
         if(arch.target().isRDNA4GPU())
         {
+            if(solution.macK == -1)
+                solution.macK = 64;
+
             if((types.typeA == "half" && types.typeB == "half")
                || (types.typeA == "bf16" && types.typeB == "bf16")
                || (types.typeA == "fp8" && types.typeB == "fp8")
@@ -1646,6 +1657,20 @@ int main(int argc, const char* argv[])
         Throw<FatalError>("Unsupported arch for GEMM client: ", arch.target().toString());
     }
 
+    if(solution.macK < solution.waveK || solution.macM < solution.waveM
+       || solution.macN < solution.waveN)
+    {
+        Throw<FatalError>(
+            fmt::format("Macro tile sizes must be greater than or equal to MI tile sizes. "
+                        "macM: {} waveM: {} macN: {} waveN: {} macK: {} waveK: {}",
+                        solution.macM,
+                        solution.waveM,
+                        solution.macN,
+                        solution.waveN,
+                        solution.macK,
+                        solution.waveK));
+    }
+
     if(types.scaleSkipPermlane)
     {
         AssertFatal(types.transA == Client::GEMMClient::TransposeType::T, ShowValue(types));

shared/rocroller/scripts/lib/rrperf/args.py

@@ -42,6 +42,12 @@ def suite(parser: argparse.ArgumentParser):
     parser.add_argument("--suite", help="Benchmark suite to run.")
 
 
+def id_filter(parser: argparse.ArgumentParser):
+    parser.add_argument(
+        "--id-filter", help="Filter benchmarks in a suite by id.", nargs="*", type=str
+    )
+
+
 def x_value(parser: argparse.ArgumentParser):
     parser.add_argument(
         "--x_value",

shared/rocroller/scripts/lib/rrperf/autoperf.py

@@ -147,6 +147,7 @@ def get_args(parser: argparse.ArgumentParser):
         args.group_results,
         args.rundir,
         args.suite,
+        args.id_filter,
     ]
     for arg in common_args:
         arg(parser)
@@ -195,7 +196,7 @@ def autoperf(
     current: bool = False,
     ancestral: bool = False,
     suite: str = None,
-    filter=None,
+    id_filter=None,
     normalize=False,
     y_zero=False,
     plot_median=False,
@@ -243,7 +244,11 @@ def autoperf(
             target,
         )
         target_success, result_dir = suite_run.run_cli(
-            build_dir=build_dir, rundir=rundir, suite=suite, filter=filter, recast=True
+            build_dir=build_dir,
+            rundir=rundir,
+            suite=suite,
+            id_filter=id_filter,
+            recast=True,
         )
         if target in no_fail_targets:
             success_no_fail &= target_success

shared/rocroller/scripts/lib/rrperf/generate.py

@@ -94,6 +94,9 @@ def generate(
     generator = rrperf.utils.empty()
     if suite is not None:
         generator = chain(generator, rrperf.utils.load_suite(suite))
+    else:
+        print("No suite specified.")
+        return
 
     if build_dir is None:
         build_dir = rrperf.utils.get_build_dir()

shared/rocroller/scripts/lib/rrperf/ls.py

@@ -30,7 +30,6 @@
 import os
 
 import rrperf.args as args
-import rrperf.rrsuites
 import rrperf.utils as utils
 
 
@@ -54,11 +53,12 @@ def run(args):
         else:
             suite = "all"
 
-    generator = getattr(rrperf.rrsuites, suite)
-    for x in generator():
+    generator = utils.load_suite(suite)
+    for x in generator:
         print(hline())
         print("-- rocRoller benchmark run")
         print("--")
+        print("Id: ", x.id)
         print("Token: ", x)
         print("")
         print(utils.sjoin(x.command()))

shared/rocroller/scripts/lib/rrperf/problems.py

@@ -26,6 +26,7 @@
 import pathlib
 from dataclasses import dataclass, field, fields, asdict
 from typing import Any, List, Optional
+from rrperf.utils import get_dataclass_id
 import yaml
 
 
@@ -180,8 +181,8 @@ class GEMMSolution:
     wave_k: int = -1
     wave_b: int = -1
 
-    workgroup_size_x: int = 64 * 2
-    workgroup_size_y: int = 2
+    workgroup_size_x: int = -1
+    workgroup_size_y: int = -1
     workgroupRemapXCC: bool = False
     workgroupRemapXCCValue: int = -1
 
@@ -234,6 +235,10 @@ class GEMM(GEMMProblem, GEMMSolution):
     def __post_init__(self):
         convert_class_params(GEMM, self)
 
+    @property
+    def id(self):
+        return get_dataclass_id(self)
+
     @property
     def run_invariant_token(self):
         return repr(GEMMProblem(**field_dict(GEMMProblem, self))) + repr(
@@ -286,7 +291,7 @@ def solution_token(self):
 class GEMMRun(GEMM):
     """GEMM run interface."""
 
-    output: pathlib.Path = field(repr=False, default=None, hash=False)
+    output: pathlib.Path = field(repr=False, default=None, hash=False, compare=False)
 
     @property
     def group(self):
@@ -407,6 +412,10 @@ class CodeGen:
     numWarmUp: int = 2
     numRuns: int = 10
 
+    @property
+    def id(self):
+        return get_dataclass_id(self)
+
     @property
     def run_invariant_token(self):
         return self.problem_token(None)
@@ -455,7 +464,7 @@ def solution_token(self):
 class CodeGenRun(CodeGen):
     """CodeGen run interface."""
 
-    output: pathlib.Path = field(repr=False, default=None, hash=False)
+    output: pathlib.Path = field(repr=False, default=None, hash=False, compare=False)
 
     @property
     def group(self):
@@ -488,8 +497,8 @@ class CodeGenResult(CodeGen, RRPerfResult):
 class TensileRun(GEMM):
     """Tensile run interface."""
 
-    config: pathlib.Path = field(repr=False, default=None, hash=False)
-    output: pathlib.Path = field(repr=False, default=None, hash=False)
+    config: pathlib.Path = field(repr=False, default=None, hash=False, compare=False)
+    output: pathlib.Path = field(repr=False, default=None, hash=False, compare=False)
     tensile_commit: str = "rocm-6.0.0"
 
     @property

shared/rocroller/scripts/lib/rrperf/run.py

@@ -62,7 +62,11 @@ def from_token(token: str):
 
 
 def run_problems(
-    generator, build_dir: Path, work_dir: Path, env: Dict[str, str]
+    generator,
+    build_dir: Path,
+    work_dir: Path,
+    env: Dict[str, str],
+    id_filter: list[str],
 ) -> bool:
 
     SOLUTION_NOT_SUPPORTED_ON_ARCH = 3
@@ -72,12 +76,13 @@ def run_problems(
     failed = []
 
     for i, problem in enumerate(generator):
-        if filter is not None:
-            pass
+        if id_filter is not None and problem.id not in id_filter:
+            continue
 
         if problem in already_run:
             continue
 
+        id = getattr(problem, "id", None)
         yaml = (work_dir / f"{problem.group}-{i:06d}.yaml").resolve()
         problem.set_output(yaml)
         cmd = problem.command()
@@ -91,6 +96,8 @@ def run_problems(
             print(f"# command: {scmd}", file=f, flush=True)
             print(f"# token: {repr(problem)}", file=f, flush=True)
             print("running:")
+            if id is not None:
+                print(f"  id: {id}")
             print(f"  command: {scmd}")
             print(f"  wrkdir:  {work_dir.resolve()}")
             print(f"  log:     {log.resolve()}")
@@ -152,6 +159,7 @@ def get_args(parser: argparse.ArgumentParser):
     common_args = [
         rrperf.args.rundir,
         rrperf.args.suite,
+        rrperf.args.id_filter,
     ]
     for arg in common_args:
         arg(parser)
@@ -163,7 +171,6 @@ def get_args(parser: argparse.ArgumentParser):
         default=False,
     )
     parser.add_argument("--token", help="Benchmark token to run.")
-    parser.add_argument("--filter", help="Filter benchmarks...")
     parser.add_argument(
         "--rocm_smi",
         default="rocm-smi",
@@ -185,7 +192,7 @@ def run_cli(
     token: str = None,
     suite: str = None,
     submit: bool = False,
-    filter: str = None,
+    id_filter: list[str] = None,
     rundir: str = None,
     build_dir: str = None,
     rocm_smi: str = "rocm-smi",
@@ -238,7 +245,7 @@ def run_cli(
     timestamp = rundir / "timestamp.txt"
     timestamp.write_text(str(datetime.datetime.now().timestamp()) + "\n")
 
-    result = run_problems(generator, build_dir, rundir, env)
+    result = run_problems(generator, build_dir, rundir, env, id_filter)
 
     if submit:
         ptsdir = rundir / "rocRoller"

shared/rocroller/scripts/lib/rrperf/utils.py

@@ -25,9 +25,12 @@
 
 import datetime
 import functools
+import json
 import os
 import subprocess
 
+from dataclasses import asdict, fields
+from hashlib import sha1
 from pathlib import Path
 
 import rrperf
@@ -99,6 +102,15 @@ def get_build_dir() -> Path:
     raise RuntimeError(f"Build directory not found.  Set {varname} to override.")
 
 
+def get_dataclass_id(obj):
+    obj_dict = asdict(obj)
+    for f in fields(obj):
+        if not f.compare:
+            del obj_dict[f.name]
+    data_str = json.dumps(obj_dict, sort_keys=True)
+    return sha1(data_str.encode()).hexdigest()
+
+
 @functools.cache
 def rocm_gfx():
     """Return GPU architecture (gfxXXXX) for local GPU device."""