Tuning quality-of-life improvements

mlir/test/perf-scripts/runtime/tuningRunner-gemm.py

@@ -1,8 +1,8 @@
 # Using a tiny GEMM with --debug-quick-tune-data. This emits a `.debug`
 # TSV of the per-config table entries (PerfConfig + TFlops) but, unlike the
-# full ``--debug`` flag, omits the heavy per-iteration ``MeasurementsMs``
-# arrays. Verify the debug file is produced, has the expected header and
-# per-config rows, and that the measurements column is absent.
+# full ``--debug`` flag, omits the per-config timing statistics in ``Stats``.
+# Verify the debug file is produced, has the expected header and per-config
+# rows, and that the statistics column is absent.
 #
 # tuningRunner.py drives real GPU tuning, so it needs the ROCm runner / GPU
 # runtime.
@@ -11,7 +11,7 @@
 # RUN: tuningRunner.py --op gemm --tuning-space=quick --debug-quick-tune-data \
 # RUN:     --config='-g 1 -m 64 -n 64 -k 64 -t f32 -out_datatype f32 -transA 0 -transB 0' \
 # RUN:     -q -o %t2.tsv
-# RUN: FileCheck %s --check-prefix=DEBUG --implicit-check-not=MeasurementsMs < %t2.tsv.debug
+# RUN: FileCheck %s --check-prefix=DEBUG --implicit-check-not=Stats < %t2.tsv.debug
 #
 # DEBUG: PerfConfig{{.*}}TFlops
 # DEBUG: {{v[0-9]+:}}

mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp

@@ -46,8 +46,10 @@
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/ErrorHandling.h"
 #include "llvm/Support/InitLLVM.h"
+#include "llvm/Support/Signals.h"
 #include "llvm/Support/SourceMgr.h"
 #include "llvm/Support/ThreadPool.h"
+#include "llvm/Support/raw_ostream.h"
 
 #include <atomic>
 #include <cassert>
@@ -57,6 +59,7 @@
 #include <mutex>
 #include <optional>
 #include <thread>
+#include <unistd.h>
 
 #include "CacheFlush.h"
 #include "ConcurrentQueue.h"
@@ -161,16 +164,16 @@ static llvm::cl::opt<unsigned> sleepUs(
 static llvm::cl::opt<bool> showStats(
     "show-stats",
     llvm::cl::desc(
-        "Print detailed stats (min, max, median, stddev, cv) in JSON format. "
-        "In case of small kernels print total_cpu_time and number of "
+        "Print detailed stats (min, max, median, stddev, cv) in nanoseconds as "
+        "JSON. In case of small kernels print total_cpu_time and number of "
         "iterations."),
     llvm::cl::init(false));
 
 static llvm::cl::opt<bool> showAllMeasurements(
     "show-all-measurements",
     llvm::cl::desc(
-        "Print all individual timing measurements in JSON format. In case of "
-        "small kernels print total_cpu_time and number of iterations."),
+        "Print all individual timing measurements in nanoseconds as JSON. In "
+        "case of small kernels print total_cpu_time and number of iterations."),
     llvm::cl::init(false));
 
 static llvm::cl::opt<std::string> benchmarkConfig(
@@ -514,7 +517,7 @@ benchmarkKernels(ArrayRef<std::string> binaries,
   // Load all modules once to reduce overhead
   std::vector<hipModule_t> modules;
   std::vector<hipFunction_t> functions;
-  auto moduleCleanup = llvm::make_scope_exit([&]() {
+  llvm::scope_exit moduleCleanup([&]() {
     for (hipModule_t mod : modules) {
       if (!mod)
         continue;
@@ -537,7 +540,7 @@ benchmarkKernels(ArrayRef<std::string> binaries,
   }
 
   // Sleep guard to avoid GPU throttling
-  auto sleepGuard = llvm::make_scope_exit([&params] {
+  llvm::scope_exit sleepGuard([&params] {
     if (params.sleepUs > 0) {
       std::this_thread::sleep_for(std::chrono::microseconds(params.sleepUs));
     }
@@ -630,9 +633,16 @@ benchmarkKernels(ArrayRef<std::string> binaries,
       return failure();
   }
 
+  constexpr auto msToNs = [](double ms) { return 1e6 * ms; };
+
+  // Convert measurements from milliseconds to nanoseconds
+  for (double &measurement : measurements) {
+    measurement = msToNs(measurement);
+  }
+
   if (params.showAllMeasurements) {
     if (isSmallKernel) {
-      llvm::outs() << "{\"total_cpu_time\":" << smallKernelCpuMs
+      llvm::outs() << "{\"total_cpu_time\":" << msToNs(smallKernelCpuMs)
                    << ",\"iterations\":" << iterations << "}\t";
     } else {
       llvm::outs() << "[";
@@ -651,7 +661,7 @@ benchmarkKernels(ArrayRef<std::string> binaries,
     // We cannot show the rest of the stats because the small kernel case uses
     // one timer only, so we cannot actually compute the min, max, etc.
     if (isSmallKernel) {
-      llvm::outs() << "{\"total_cpu_time\":" << smallKernelCpuMs
+      llvm::outs() << "{\"total_cpu_time\":" << msToNs(smallKernelCpuMs)
                    << ",\"iterations\":" << iterations << "}\t";
     }
     if (measurements.size() > 1) {
@@ -667,11 +677,10 @@ benchmarkKernels(ArrayRef<std::string> binaries,
     }
   }
 
-  auto msToNs = [](double ms) { return 1e6 * ms; };
   if (params.useMedian)
-    return msToNs(computeMedian(measurements));
+    return computeMedian(measurements);
   else
-    return msToNs(computeMean(trimValues(measurements, params.trimPercent)));
+    return computeMean(trimValues(measurements, params.trimPercent));
 }
 
 static int toKernelOrder(Attribute attr) {
@@ -702,6 +711,46 @@ static LogicalResult extractFuncOps(ModuleOp op,
   return success();
 }
 
+/// Perf config the calling thread is compiling, or null when it is between
+/// configs. Points into the config list, which outlives the workers.
+static thread_local const SmallString<64> *compilingConfig = nullptr;
+
+/// Writes to stderr without llvm::errs(), which is neither async-signal-safe
+/// nor safe to reenter from a fatal error handler. Give up silently on a failed
+/// write: there is nothing left to report it with.
+static void writeToStderr(StringRef message) {
+  ssize_t written = ::write(STDERR_FILENO, message.data(), message.size());
+  (void)written;
+}
+
+/// Names the perf config being compiled on the crashing thread. Clears the
+/// pointer so that the abort() following report_fatal_error() does not report
+/// the same config twice through the signal handler.
+static void reportCompilingConfig() {
+  if (!compilingConfig)
+    return;
+  SmallString<256> message;
+  {
+    llvm::raw_svector_ostream os(message);
+    os << "Offending perf config: " << *compilingConfig << "\n"
+       << "Reproduce with `--benchmark-config=" << *compilingConfig << "`\n";
+  }
+  writeToStderr(message);
+  compilingConfig = nullptr;
+}
+
+static void compilationFatalErrorHandler(void *, const char *reason, bool) {
+  SmallString<256> message;
+  {
+    llvm::raw_svector_ostream os(message);
+    os << "LLVM ERROR: " << reason << "\n";
+  }
+  writeToStderr(message);
+  reportCompilingConfig();
+}
+
+static void compilationSignalHandler(void *) { reportCompilingConfig(); }
+
 static bool doesModuleHaveFusions(ModuleOp module) {
   WalkResult result = module.walk([](Operation *op) {
     // Check for linalg.generic or rock.reduce (standalone fusion ops)
@@ -767,7 +816,7 @@ static LogicalResult runTuningLoop(ModuleOp source) {
   // 3. Create HIP stream and allocate device buffers
   hipStream_t stream;
   HIPCHECK(hipStreamCreate(&stream));
-  auto streamCleanup = llvm::make_scope_exit([&]() {
+  llvm::scope_exit streamCleanup([&]() {
     hipError_t status = hipStreamDestroy(stream);
     if (status != hipSuccess) {
       llvm::errs() << "HIP error in hipStreamDestroy: "
@@ -776,7 +825,7 @@ static LogicalResult runTuningLoop(ModuleOp source) {
   });
 
   std::vector<void *> gpuBuffers;
-  auto bufferCleanup = llvm::make_scope_exit([&]() {
+  llvm::scope_exit bufferCleanup([&]() {
     for (void *buffer : gpuBuffers) {
       // hipFree does not allow nullptrs, so make sure to check for it first
       if (!buffer)
@@ -1020,7 +1069,11 @@ static LogicalResult runTuningLoop(ModuleOp source) {
         if (idx >= configs.size())
           break;
 
-        if (!compilationResults.push(compileConfig(idx, myRes)))
+        compilingConfig = &configs[idx];
+        CompilationResult result = compileConfig(idx, myRes);
+        compilingConfig = nullptr;
+
+        if (!compilationResults.push(std::move(result)))
           break; // Queue terminated
       }
 
@@ -1036,7 +1089,7 @@ static LogicalResult runTuningLoop(ModuleOp source) {
       threads.emplace_back(worker);
     }
 
-    auto threadCleanup = llvm::make_scope_exit([&] {
+    llvm::scope_exit threadCleanup([&] {
       // In case of early termination, signal all threads to stop
       compilationResults.terminate();
       for (auto &t : threads) {
@@ -1104,6 +1157,11 @@ static LogicalResult runTuningLoop(ModuleOp source) {
 int main(int argc, char **argv) {
   llvm::InitLLVM y(argc, argv);
 
+  // Name the perf config under compilation if a backend pass in the pipeline
+  // dies, either through report_fatal_error or a crash signal.
+  llvm::install_fatal_error_handler(compilationFatalErrorHandler);
+  llvm::sys::AddSignalHandler(compilationSignalHandler, nullptr);
+
   mlir::registerMLIRCLOptions();
   llvm::cl::ParseCommandLineOptions(argc, argv, "rocMLIR tuning driver");
 

mlir/utils/performance/perfRunner.py

@@ -259,6 +259,16 @@ def chip_has_mfma():
     return has_feature(lookup_arch_info(get_chip()).default_features, GemmFeatures.MFMA)
 
 
+def chip_uses_wgp_mode(chip: Optional[str] = None) -> bool:
+    """Whether the chip is built in workgroup processor (WGP) mode.
+
+    In WGP mode HIP counts one multiprocessor per two-CU workgroup processor, so the CU
+    count it reports is half of the physical one that rocminfo prints. Wave size is the
+    same proxy for WGP mode that fixNaviProperties() uses in AmdArchDb.cpp.
+    """
+    return lookup_arch_info(chip or get_chip()).wave_size == 32
+
+
 DATA_TYPES_ATTENTION = None
 
 

mlir/utils/performance/tests/test_perfRunner.py

@@ -168,6 +168,31 @@ def fake_lookup(arch):
         assert captured["arch"] == "native:1"
 
 
+class TestChipUsesWgpMode:
+    """Tests for chip_uses_wgp_mode (wave size proxy for WGP mode)."""
+
+    def test_wave64_chip_is_not_wgp_mode(self, monkeypatch):
+        monkeypatch.setattr(perfRunner, "lookup_arch_info",
+                            lambda arch: types.SimpleNamespace(wave_size=64))
+        assert not perfRunner.chip_uses_wgp_mode("gfx942")
+
+    def test_wave32_chip_is_wgp_mode(self, monkeypatch):
+        monkeypatch.setattr(perfRunner, "lookup_arch_info",
+                            lambda arch: types.SimpleNamespace(wave_size=32))
+        assert perfRunner.chip_uses_wgp_mode("gfx1100")
+
+    def test_defaults_to_current_chip(self, monkeypatch):
+        captured = {}
+
+        def fake_lookup(arch):
+            captured["arch"] = arch
+            return types.SimpleNamespace(wave_size=32)
+
+        monkeypatch.setattr(perfRunner, "lookup_arch_info", fake_lookup)
+        assert perfRunner.chip_uses_wgp_mode()
+        assert captured["arch"] == "gfx900"
+
+
 class TestParseDataTypes:
     """Tests for parse_data_types (gemm data types)."""