[rocRoller] Add unique IDs to rrsuite

shared/rocroller/client/src/codegen_stress_test.cpp

@@ -116,7 +116,7 @@ Generator<Instruction> simple_mi(ContextPtr m_context)
     {
         mi_mnemonic = "v_mfma_f32_32x32x1f32";
     }
-    else if(arch.HasCapability(GPUCapability::HasWMMA) && arch.target().isRDNA4GPU())
+    else if(arch.HasCapability(GPUCapability::HasWMMA_f32_16x16x16_f16))
     {
         mi_mnemonic = "v_wmma_f32_16x16x16_f16";
     }
@@ -140,7 +140,7 @@ Generator<Instruction> complex_mi_with_coop(ContextPtr m_context)
     {
         mi_mnemonic = "v_mfma_f32_32x32x1f32";
     }
-    else if(arch.HasCapability(GPUCapability::HasWMMA) && arch.target().isRDNA4GPU())
+    else if(arch.HasCapability(GPUCapability::HasWMMA_f32_16x16x16_f16))
     {
         mi_mnemonic = "v_wmma_f32_16x16x16_f16";
     }

shared/rocroller/client/src/gemm.cpp

@@ -1081,14 +1081,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,
@@ -1278,7 +1278,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.");
@@ -1539,8 +1539,19 @@ int main(int argc, const char* argv[])
         types.scaleBlockSize = arch.GetCapability(GPUCapability::DefaultScaleBlockSize);
     }
 
-    AssertFatal((types.typeAcc == "float") || (types.typeAcc == "half")
-                || (types.typeAcc == "bf16"));
+    if(solution.workgroupSizeX == -1)
+        solution.workgroupSizeX = 2 * arch.GetCapability(GPUCapability::DefaultWavefrontSize);
+    if(solution.workgroupSizeY == -1)
+        solution.workgroupSizeY = 2;
+
+    const DataType typeA   = fromString<DataType>(types.typeA);
+    const DataType typeB   = fromString<DataType>(types.typeB);
+    const DataType typeC   = fromString<DataType>(types.typeC);
+    const DataType typeD   = fromString<DataType>(types.typeD);
+    const DataType typeAcc = fromString<DataType>(types.typeAcc);
+
+    AssertFatal((typeAcc == DataType::Float) || (typeAcc == DataType::Half)
+                || (typeAcc == DataType::BFloat16));
 
     // TODO: Reevaluate the relationship between problem and solution params.
     problem.workgroupMappingDim = solution.workgroupMappingDim;
@@ -1550,114 +1561,138 @@ 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(types.typeA == "float" && types.typeB == "float" && types.typeC == "float"
-           && types.typeD == "float")
+        if(solution.macK == -1)
+            solution.macK = 64;
+        if(solution.waveB == -1)
+            solution.waveB = 1;
+
+        if(typeA == DataType::Float && typeB == DataType::Float && typeC == DataType::Float
+           && typeD == DataType::Float)
         {
             if(solution.waveM == -1)
                 solution.waveM = 32;
             if(solution.waveN == -1)
                 solution.waveN = 32;
             if(solution.waveK == -1)
                 solution.waveK = 2;
-            if(solution.waveB == -1)
-                solution.waveB = 1;
         }
-        else if(types.typeA == "half" && types.typeB == "half")
+        else if(typeA == DataType::Half && typeB == DataType::Half)
         {
             if(solution.waveM == -1)
                 solution.waveM = 32;
             if(solution.waveN == -1)
                 solution.waveN = 32;
             if(solution.waveK == -1)
                 solution.waveK = 8;
-            if(solution.waveB == -1)
-                solution.waveB = 1;
         }
-        else if(types.typeA == "bf16" && types.typeB == "bf16")
+        else if(typeA == DataType::BFloat16 && typeB == DataType::BFloat16)
         {
             if(solution.waveM == -1)
                 solution.waveM = 16;
             if(solution.waveN == -1)
                 solution.waveN = 16;
             if(solution.waveK == -1)
                 solution.waveK = 8;
-            if(solution.waveB == -1)
-                solution.waveB = 1;
         }
-        else if((types.typeA == "fp8" && types.typeB == "fp8")
-                || (types.typeA == "bf8" && types.typeB == "bf8"))
+        else if((typeA == DataType::FP8 && typeB == DataType::FP8)
+                || (typeA == DataType::BF8 && typeB == DataType::BF8))
         {
             if(solution.waveM == -1)
                 solution.waveM = 16;
             if(solution.waveN == -1)
                 solution.waveN = 16;
             if(solution.waveK == -1)
                 solution.waveK = 32;
-            if(solution.waveB == -1)
-                solution.waveB = 1;
         }
     }
     else if(arch.HasCapability(GPUCapability::HasWMMA))
     {
-        if(arch.target().isRDNA4GPU())
-        {
-            if((types.typeA == "half" && types.typeB == "half")
-               || (types.typeA == "bf16" && types.typeB == "bf16")
-               || (types.typeA == "fp8" && types.typeB == "fp8")
-               || (types.typeA == "bf8" && types.typeB == "bf8")
-               || (types.typeA == "bf8" && types.typeB == "fp8")
-               || (types.typeA == "fp8" && types.typeB == "bf8"))
+        if(solution.waveM == -1)
+            solution.waveM = 16;
+        if(solution.waveN == -1)
+            solution.waveN = 16;
+        if(solution.waveB == -1)
+            solution.waveB = 1;
+
+        if((typeA == DataType::Half && typeB == DataType::Half)
+           || (typeA == DataType::BFloat16 && typeB == DataType::BFloat16))
+        {
+            if(solution.macK == -1)
+                solution.macK = 64;
+
+            if(arch.HasCapability(GPUCapability::HasWMMA_f32_16x16x16_f16))
             {
-                if(solution.waveM == -1)
-                    solution.waveM = 16;
-                if(solution.waveN == -1)
-                    solution.waveN = 16;
                 if(solution.waveK == -1)
                     solution.waveK = 16;
-                if(solution.waveB == -1)
-                    solution.waveB = 1;
             }
-            else
+        }
+        else if(isUnpackedF8(fromString<DataType>(solution.types.typeA))
+                && isUnpackedF8(fromString<DataType>(solution.types.typeB)))
+        {
+            if(solution.macK == -1)
+                solution.macK = 64;
+
+            if(arch.HasCapability(GPUCapability::HasWMMA_f32_16x16x16_f8))
             {
-                // Override default settings for the `example` and `generate` subcommands.
-                if(example->parsed() || generate->parsed())
-                {
-                    types.typeA    = "half";
-                    types.typeB    = "half";
-                    types.typeC    = "half";
-                    types.typeD    = "half";
-                    solution.waveM = 16;
-                    solution.waveN = 16;
+                if(solution.waveK == -1)
                     solution.waveK = 16;
-                    solution.waveB = 1;
-                }
-                else
-                {
-                    Throw<FatalError>("Unsupported MI on: ",
-                                      arch.target().toString(),
-                                      ShowValue(types.typeA),
-                                      ShowValue(types.typeB),
-                                      ShowValue(types.typeC),
-                                      ShowValue(types.typeD),
-                                      ShowValue(types.typeAcc));
-                }
             }
-            // TODO Support prefetch on gfx12
-            solution.prefetch = false;
-        }
-        else
-        {
-            Throw<FatalError>("Unsupported arch for GEMM client: ", arch.target().toString());
         }
     }
     else
     {
         Throw<FatalError>("Unsupported arch for GEMM client: ", arch.target().toString());
     }
 
+    if(arch.target().isRDNA4GPU())
+    {
+        // Override default settings for the `example` and `generate` subcommands.
+        if((example->parsed() || generate->parsed()) && typeA == DataType::Float
+           && typeB == DataType::Float)
+        {
+            std::cout << "Warning: A and B types and wave sizes have been overridden for RDNA4."
+                      << std::endl;
+            types.typeA    = "half";
+            types.typeB    = "half";
+            types.typeC    = "half";
+            types.typeD    = "half";
+            solution.waveM = 16;
+            solution.waveN = 16;
+            solution.waveK = 16;
+            solution.waveB = 1;
+            solution.macK  = 64;
+        }
+
+        if(solution.prefetch)
+        {
+            std::cout << "Warning: disabling prefetching for RDNA4." << std::endl;
+            solution.prefetch = false;
+        }
+    }
+
+    AssertFatal(solution.waveM > 0 && solution.waveN > 0 && solution.waveK > 0
+                    && solution.waveB > 0,
+                fmt::format("MI tile sizes must be set greater than zero. "
+                            "waveM: {} waveN: {} waveK: {} waveB: {}",
+                            solution.waveM,
+                            solution.waveN,
+                            solution.waveK,
+                            solution.waveB));
+
+    AssertFatal(solution.macK >= solution.waveK && solution.macM >= solution.waveM
+                    && solution.macN >= solution.waveN,
+                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

@@ -37,16 +37,27 @@
 
 
 def generate_kernels(
-    generator, architecture: str, build_dir: Path, work_dir: Path, env: Dict[str, str]
+    generator,
+    architecture: str,
+    build_dir: Path,
+    work_dir: Path,
+    env: Dict[str, str],
+    id_filter: list[str],
 ) -> bool:
 
     already_run = set()
     failed = []
 
     for i, problem in enumerate(generator):
+        if id_filter is not None and not any(
+            problem.id.startswith(filt) for filt in id_filter
+        ):
+            continue
+
         if problem in already_run:
             continue
 
+        id = getattr(problem, "id", None)
         cmd = problem.command(architecture=architecture, generate_only=True)
         log = (work_dir / f"{problem.group}-{i:06d}.log").resolve()
         rr_env = {k: str(v) for k, v in env.items() if k.startswith("ROC")}
@@ -57,6 +68,7 @@ def generate_kernels(
             print(f"# command: {sjoin(cmd)}", file=f, flush=True)
             print(f"# token: {repr(problem)}", file=f, flush=True)
             print("running:")
+            print(f"  id: {id}")
             print(f"  command: {sjoin(cmd)}")
             print(f"  wrkdir:  {work_dir.resolve()}")
             print(f"  log:     {log.resolve()}")
@@ -72,6 +84,11 @@ def generate_kernels(
         already_run.add(problem)
 
     if len(failed) > 0:
+        ids = [getattr(problem, "id", None) for i, problem in failed]
+        print("")
+        print(f"Failed {len(failed)} problems ids:")
+        print(" ".join([str(id) for id in ids]))
+        print("")
         print(f"Failed {len(failed)} problems:")
         for i, problem in failed:
             cmd = list(map(str, problem.command()))
@@ -85,13 +102,17 @@ def generate(
     suite: str = None,
     rundir: str = None,
     build_dir: str = None,
+    id_filter: list[str] = None,
     **kwargs,
 ):
     """Generate kernels!"""
 
     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()
@@ -111,7 +132,9 @@ def generate(
     except Exception:
         git_commit.write_text("NO_COMMIT\n")
 
-    success = generate_kernels(generator, architecture, build_dir, run_dir, env)
+    success = generate_kernels(
+        generator, architecture, build_dir, run_dir, env, id_filter
+    )
 
     if not success:
         raise RuntimeError("Some jobs failed.")
@@ -121,6 +144,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)