From 0c355e4d34ce43f583e867cdcec1dd4e0db2e2eb Mon Sep 17 00:00:00 2001 From: Daniel Hernandez Date: Thu, 11 Jun 2026 10:03:16 +0200 Subject: [PATCH 1/6] Align rocmlir-tuning-driver benchmarking with Triton do_bench --- .../rocmlir-tuning-driver.cpp | 313 ++++++++---------- mlir/utils/performance/perfRunner.py | 13 +- mlir/utils/performance/tuningRunner.py | 12 +- 3 files changed, 160 insertions(+), 178 deletions(-) diff --git a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp index 9e0111a5a2a1..50cff03b2a31 100644 --- a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp +++ b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp @@ -133,14 +133,19 @@ static llvm::cl::opt tuningSpaceKind( llvm::cl::value_desc("tuning space to use"), llvm::cl::init(rock::TuningParamSetKind::Full)); -static llvm::cl::opt numIterations( - "num-iterations", - llvm::cl::desc("Number of times to run each kernel for averaging"), - llvm::cl::value_desc("number of runs"), llvm::cl::init(100)); - -static llvm::cl::opt warmupIterations( - "warmup-iterations", llvm::cl::desc("Number of warmup runs"), - llvm::cl::value_desc("number of warmup runs"), llvm::cl::init(10)); +static llvm::cl::opt rep( + "rep", + llvm::cl::desc("Target benchmark time in milliseconds. The number of " + "measured iterations is derived from this budget and the " + "estimated per-launch runtime (Triton do_bench style)."), + llvm::cl::value_desc("benchmark milliseconds"), llvm::cl::init(100)); + +static llvm::cl::opt warmup( + "warmup", + llvm::cl::desc("Target warmup time in milliseconds. The number of warmup " + "iterations is derived from this budget and the estimated " + "per-launch runtime (Triton do_bench style)."), + llvm::cl::value_desc("warmup milliseconds"), llvm::cl::init(25)); static llvm::cl::opt useMedian("use-median", @@ -161,16 +166,12 @@ static llvm::cl::opt sleepUs( static llvm::cl::opt 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 " - "iterations."), + "Print detailed stats (min, max, median, stddev, cv) in JSON format."), llvm::cl::init(false)); static llvm::cl::opt 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."), + llvm::cl::desc("Print all individual timing measurements in JSON format."), llvm::cl::init(false)); static llvm::cl::opt benchmarkConfig( @@ -328,8 +329,8 @@ static std::vector trimValues(const std::vector &values, } struct BenchmarkParams { - unsigned numIterations; - unsigned warmupIterations; + unsigned warmupMs; + unsigned repMs; bool useMedian; unsigned trimPercent; unsigned sleepUs; @@ -409,54 +410,38 @@ struct ThreadResources { bool isValid() const { return sourceModule && *sourceModule; } }; -static LogicalResult measureSmallKernel( - unsigned iterations, hipStream_t stream, - const std::vector &functions, ArrayRef blockSizes, - ArrayRef gridSizes, std::vector &argPointers, - std::vector &measurements, double &smallKernelCpuMs, - bool benchmarkMode, const std::optional &gpuRunDeadline, - unsigned timeoutSec, StringRef perfConfig) { - // Special case for small kernels, where we measure the time for all kernels - // at once, using CPU timers. - auto iterationStart = std::chrono::steady_clock::now(); - for (unsigned iter = 0; iter < iterations; ++iter) { - // Do not flush caches in benchmark mode, as we do not want to - // time the cache flush (it's okay if we are running in tuning mode). - if (!benchmarkMode) { - if (failed(flushInstructionCache(stream))) { - return failure(); - } - if (failed(flushL2Cache(stream))) { - return failure(); - } +static LogicalResult +measureKernel(unsigned iterations, hipStream_t stream, + const std::vector &functions, + ArrayRef blockSizes, ArrayRef gridSizes, + std::vector &argPointers, + std::vector &measurements, + const std::optional &gpuRunDeadline, + unsigned timeoutSec, StringRef perfConfig) { + // Pre-allocate one event pair per iteration so we can record them all in a + // tight loop and synchronize only once at the end. This matches Triton's + // do_bench, which minimizes host-side overhead between launches (no + // per-iteration synchronization). + std::vector startEvents(iterations, nullptr); + std::vector stopEvents(iterations, nullptr); + auto eventCleanup = llvm::make_scope_exit([&]() { + for (hipEvent_t event : startEvents) { + if (event) + (void)hipEventDestroy(event); } - for (auto [func, blockSize, gridSize] : - llvm::zip(functions, blockSizes, gridSizes)) { - HIPCHECK(hipExtModuleLaunchKernel( - func, gridSize * blockSize, 1, 1, blockSize, 1, 1, 0, stream, - argPointers.data(), nullptr, nullptr, nullptr)); + for (hipEvent_t event : stopEvents) { + if (event) + (void)hipEventDestroy(event); } + }); + for (unsigned iter = 0; iter < iterations; ++iter) { + HIPCHECK(hipEventCreate(&startEvents[iter])); + HIPCHECK(hipEventCreate(&stopEvents[iter])); } - if (failed(synchronizeStreamWithTimeout(stream, gpuRunDeadline, timeoutSec, - perfConfig, "measurement"))) - return failure(); - smallKernelCpuMs = std::chrono::duration( - std::chrono::steady_clock::now() - iterationStart) - .count(); - measurements.push_back(smallKernelCpuMs / iterations); - return success(); -} - -static LogicalResult -measureLargeKernel(unsigned iterations, hipStream_t stream, - const std::vector &functions, - ArrayRef blockSizes, ArrayRef gridSizes, - std::vector &argPointers, - std::vector &measurements, - const std::optional &gpuRunDeadline, - unsigned timeoutSec, StringRef perfConfig) { - // Measure runs normally. + // Record all iterations back-to-back. Each measurement brackets the full + // kernel chain (one start before, one stop after), matching how Triton times + // the whole callable rather than each kernel individually. for (unsigned iter = 0; iter < iterations; ++iter) { if (failed(flushInstructionCache(stream))) { return failure(); @@ -465,32 +450,26 @@ measureLargeKernel(unsigned iterations, hipStream_t stream, return failure(); } - double totalMilliseconds = 0.0; - + HIPCHECK(hipEventRecord(startEvents[iter], stream)); for (auto [func, blockSize, gridSize] : llvm::zip(functions, blockSizes, gridSizes)) { - hipEvent_t startEvent, stopEvent; - HIPCHECK(hipEventCreate(&startEvent)); - HIPCHECK(hipEventCreate(&stopEvent)); - HIPCHECK(hipExtModuleLaunchKernel( func, gridSize * blockSize, 1, 1, blockSize, 1, 1, 0, stream, - argPointers.data(), nullptr, startEvent, stopEvent)); - if (failed(synchronizeStreamWithTimeout( - stream, gpuRunDeadline, timeoutSec, perfConfig, "measurement"))) - return failure(); - - float currentMilliseconds = 0.0; - HIPCHECK( - hipEventElapsedTime(¤tMilliseconds, startEvent, stopEvent)); - - HIPCHECK(hipEventDestroy(stopEvent)); - HIPCHECK(hipEventDestroy(startEvent)); - - totalMilliseconds += static_cast(currentMilliseconds); + argPointers.data(), nullptr, nullptr, nullptr)); } + HIPCHECK(hipEventRecord(stopEvents[iter], stream)); + } + + // Single synchronization after all iterations have been queued. + if (failed(synchronizeStreamWithTimeout(stream, gpuRunDeadline, timeoutSec, + perfConfig, "measurement"))) + return failure(); - measurements.push_back(totalMilliseconds); + for (unsigned iter = 0; iter < iterations; ++iter) { + float currentMilliseconds = 0.0; + HIPCHECK(hipEventElapsedTime(¤tMilliseconds, startEvents[iter], + stopEvents[iter])); + measurements.push_back(static_cast(currentMilliseconds)); } return success(); @@ -503,8 +482,6 @@ benchmarkKernels(ArrayRef binaries, ArrayRef gridSizes, MutableArrayRef gpuBuffers, hipStream_t stream, StringRef perfConfig, const BenchmarkParams ¶ms) { - bool benchmarkMode = !params.benchmarkConfig.empty(); - // HIP wants an array of pointers to each argument std::vector argPointers; for (void *&item : gpuBuffers) { @@ -549,111 +526,95 @@ benchmarkKernels(ArrayRef binaries, std::optional gpuRunDeadline = makeTimeoutDeadline(params.gpuRunTimeoutSec); - bool isSmallKernel = false; - unsigned iterations = params.numIterations; - - if (params.warmupIterations > 0) { - // Warmup run. We measure the warmup to get an estimate of the kernel - // runtime. We will use this estimate to determine if the kernel is small or - // not. - double totalMillisecondsWarmup = 0.0; - for (unsigned iter = 0; iter < params.warmupIterations; ++iter) { + // Estimate the per-launch runtime so we can size warmup/benchmark iteration + // counts from the requested time budgets (Triton do_bench style). We time a + // handful of launches (flushing caches between them) using a single event + // pair. + constexpr unsigned estimateRuns = 5; + double estimateMs = 0.0; + { + hipEvent_t startEvent, stopEvent; + HIPCHECK(hipEventCreate(&startEvent)); + HIPCHECK(hipEventCreate(&stopEvent)); + + HIPCHECK(hipEventRecord(startEvent, stream)); + for (unsigned iter = 0; iter < estimateRuns; ++iter) { + // The cache flushes are inside the timed window here (unlike the actual + // measurement loop, which flushes before recording the start event). This + // is intentional, to match Triton's do_bench, whose estimate loop also + // includes the cache clear. + if (failed(flushInstructionCache(stream))) + return failure(); + if (failed(flushL2Cache(stream))) + return failure(); for (auto [func, blockSize, gridSize] : llvm::zip(functions, blockSizes, gridSizes)) { - hipEvent_t startEvent, stopEvent; - HIPCHECK(hipEventCreate(&startEvent)); - HIPCHECK(hipEventCreate(&stopEvent)); - HIPCHECK(hipExtModuleLaunchKernel( func, gridSize * blockSize, 1, 1, blockSize, 1, 1, 0, stream, - argPointers.data(), nullptr, startEvent, stopEvent)); - - if (failed(synchronizeStreamWithTimeout(stream, gpuRunDeadline, - params.gpuRunTimeoutSec, - perfConfig, "warmup"))) - return failure(); - - float currentMilliseconds = 0.0; - HIPCHECK( - hipEventElapsedTime(¤tMilliseconds, startEvent, stopEvent)); - - HIPCHECK(hipEventDestroy(stopEvent)); - HIPCHECK(hipEventDestroy(startEvent)); - - // hipEventElapsedTime seemingly can return negative values for fast - // kernels due to GPU clock precision issues. This is extremely relevant - // when we have a small number of warmup iterations (e.g., 1) for small - // kernels. Clamp to the documented resolution of ~1 microsecond - // (0.001 ms) if this is the case. - if (currentMilliseconds < 0.0f) { - constexpr float minMeasurableMs = 0.001f; - currentMilliseconds = minMeasurableMs; - } - - totalMillisecondsWarmup += static_cast(currentMilliseconds); + argPointers.data(), nullptr, nullptr, nullptr)); } } - totalMillisecondsWarmup /= params.warmupIterations; - assert(totalMillisecondsWarmup >= 0.0f && - "totalMillisecondsWarmup must be greater than 0"); - - // We want to get at least 1ms of kernel execution time - // (counting all iterations), so increase the number of iterations - // if necessary. - constexpr float minTotalMilliseconds = 1.0f; - iterations = std::max( - iterations, static_cast(std::ceil(minTotalMilliseconds / - totalMillisecondsWarmup))); - - // Depending on the runtime of the kernel, - // we will use a different approach to measure the runs. - // We consider a kernel to be small if a single iteration takes less than - // 1ms to run. - constexpr float smallKernelThreshold = 1.0f; - isSmallKernel = totalMillisecondsWarmup < smallKernelThreshold; + HIPCHECK(hipEventRecord(stopEvent, stream)); + if (failed(synchronizeStreamWithTimeout(stream, gpuRunDeadline, + params.gpuRunTimeoutSec, perfConfig, + "warmup"))) + return failure(); + + float elapsedMs = 0.0; + HIPCHECK(hipEventElapsedTime(&elapsedMs, startEvent, stopEvent)); + HIPCHECK(hipEventDestroy(stopEvent)); + HIPCHECK(hipEventDestroy(startEvent)); + + estimateMs = static_cast(elapsedMs) / estimateRuns; + // hipEventElapsedTime can return tiny/negative values for very fast kernels + // due to GPU clock precision. Clamp to the documented ~1 microsecond + // resolution (0.001 ms) to avoid divide-by-zero / overflow below. + constexpr double minMeasurableMs = 0.001; + if (estimateMs < minMeasurableMs) + estimateMs = minMeasurableMs; } + // Derive iteration counts from the time budgets, like Triton's do_bench. + unsigned nWarmup = std::max( + 1, static_cast(params.warmupMs / estimateMs)); + unsigned iterations = + std::max(1, static_cast(params.repMs / estimateMs)); + + // Warm-up (untimed): just run the kernel chain nWarmup times. + for (unsigned iter = 0; iter < nWarmup; ++iter) { + for (auto [func, blockSize, gridSize] : + llvm::zip(functions, blockSizes, gridSizes)) { + HIPCHECK(hipExtModuleLaunchKernel( + func, gridSize * blockSize, 1, 1, blockSize, 1, 1, 0, stream, + argPointers.data(), nullptr, nullptr, nullptr)); + } + } + if (failed(synchronizeStreamWithTimeout(stream, gpuRunDeadline, + params.gpuRunTimeoutSec, perfConfig, + "warmup"))) + return failure(); + // Measure runs std::vector measurements; - double smallKernelCpuMs = 0.0; - if (isSmallKernel) { - if (failed(measureSmallKernel( - iterations, stream, functions, blockSizes, gridSizes, argPointers, - measurements, smallKernelCpuMs, benchmarkMode, gpuRunDeadline, - params.gpuRunTimeoutSec, perfConfig))) - return failure(); - } else { - if (failed(measureLargeKernel( - iterations, stream, functions, blockSizes, gridSizes, argPointers, - measurements, gpuRunDeadline, params.gpuRunTimeoutSec, perfConfig))) - return failure(); - } + if (failed(measureKernel(iterations, stream, functions, blockSizes, gridSizes, + argPointers, measurements, gpuRunDeadline, + params.gpuRunTimeoutSec, perfConfig))) + return failure(); if (params.showAllMeasurements) { - if (isSmallKernel) { - llvm::outs() << "{\"total_cpu_time\":" << smallKernelCpuMs - << ",\"iterations\":" << iterations << "}\t"; - } else { - llvm::outs() << "["; - for (size_t i = 0; i < measurements.size(); ++i) { - if (i > 0) - llvm::outs() << ","; - llvm::outs() << measurements[i]; - } - llvm::outs() << "]\t"; + llvm::outs() << "["; + for (size_t i = 0; i < measurements.size(); ++i) { + if (i > 0) + llvm::outs() << ","; + llvm::outs() << measurements[i]; } + llvm::outs() << "]\t"; } std::sort(measurements.begin(), measurements.end()); if (params.showStats) { - // 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 - << ",\"iterations\":" << iterations << "}\t"; - } if (measurements.size() > 1) { float median = computeMedian(measurements); float min = measurements.front(); @@ -810,10 +771,18 @@ static LogicalResult runTuningLoop(ModuleOp source) { // NOTE: Compilation (PassManager::run()) resets the cl opts, so we have to // save the values. - const BenchmarkParams benchmarkParams = { - numIterations, warmupIterations, useMedian, trimPercent, - sleepUs, showStats, showAllMeasurements, tuningSpaceKind, - numCompileThreads, benchmarkConfig, waitForCompiles, gpuRunTimeout}; + const BenchmarkParams benchmarkParams = {warmup, + rep, + useMedian, + trimPercent, + sleepUs, + showStats, + showAllMeasurements, + tuningSpaceKind, + numCompileThreads, + benchmarkConfig, + waitForCompiles, + gpuRunTimeout}; rock::TuningParamSetKind effectiveKind = benchmarkParams.tuningSpaceKind; unsigned numTuningIterations = rock::getNumberOfIterations(effectiveKind); diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py index 060f87b33480..dc602cfde6f2 100644 --- a/mlir/utils/performance/perfRunner.py +++ b/mlir/utils/performance/perfRunner.py @@ -57,8 +57,15 @@ 'bf8_bf8': 'f32', 'f4E2M1FN': 'f32' } +# rocmlir-gen host-harness kernel repeat count (--kernel-repeats, used with -ph). MLIR_N_REPEATS = 100 -WARMUP_ITERATIONS = 10 + +# Time budgets (ms) for the tuning-driver benchmark. The number of warmup and +# measured iterations is derived from these budgets and the estimated per-launch +# runtime (Triton do_bench style). Deliberately stricter than Triton's defaults +# (25/100) to get more stable performance numbers. +BENCH_WARMUP_MS = 50 +BENCH_REP_MS = 200 SLEEP_US = 1000 # 1 ms FILTER_LAYOUT_MAP = {'N': 'k', 'C': 'c', 'H': 'y', 'W': 'x', 'G': 'g', '0': '0', '1': '1'} @@ -2028,8 +2035,8 @@ def run_config_with_mlir(config: PerfConfiguration, print("Using HIP timing for benchmarking") tuning_driver_command = [ paths.mlir_paths.rocmlir_tuning_driver_path, f'--benchmark-config={config.perfconfig}', - f'--num-iterations={MLIR_N_REPEATS}', f'--warmup-iterations={WARMUP_ITERATIONS}', - f'--sleep-us={SLEEP_US}', '--use-median', '-' + f'--rep={BENCH_REP_MS}', f'--warmup={BENCH_WARMUP_MS}', f'--sleep-us={SLEEP_US}', + '--use-median', '-' ] outs, noerr = run_pipeline([rocmlir_gen_cmd.split(), tuning_driver_command]) if noerr: diff --git a/mlir/utils/performance/tuningRunner.py b/mlir/utils/performance/tuningRunner.py index 7b91e6071ba3..6326dfc48429 100755 --- a/mlir/utils/performance/tuningRunner.py +++ b/mlir/utils/performance/tuningRunner.py @@ -66,8 +66,14 @@ # Constants # ============================================================================= +# rocmlir-gen host-harness kernel repeat count (--kernel-repeats, used with -ph). MLIR_N_REPEATS = 10 -WARMUP_ITERATIONS = 1 + +# Time budgets (ms) for the tuning-driver benchmark. The number of warmup and +# measured iterations is derived from these budgets and the estimated per-launch +# runtime (Triton do_bench style). These mirror Triton's do_bench defaults. +TUNE_WARMUP_MS = 25 +TUNE_REP_MS = 100 SLEEP_US = 100 # 0.1 ms # A GPU run timeout is different from the outer tuning subprocess timeout: an @@ -1450,8 +1456,8 @@ def tune_config(test_vector: str, conf_class: type, paths: Paths, options: Optio tuning_driver_args = [ f"--tuning-space={options.tuning_space_kind}", - f"--num-iterations={MLIR_N_REPEATS}", - f"--warmup-iterations={WARMUP_ITERATIONS}", + f"--rep={TUNE_REP_MS}", + f"--warmup={TUNE_WARMUP_MS}", "--use-median", f"--sleep-us={SLEEP_US}", f"--show-all-measurements={options.debug}", From a6f1363e8b9a3be189508d26de4b2a45a9db60ba Mon Sep 17 00:00:00 2001 From: Daniel Hernandez Date: Tue, 16 Jun 2026 14:08:28 +0200 Subject: [PATCH 2/6] port --flush-last-level-cache --- mlir/include/mlir/Dialect/Rock/IR/AmdArchDb.h | 4 ++ mlir/lib/Dialect/Rock/IR/AmdArchDb.cpp | 41 +++++++++++ .../rocmlir-tuning-driver/CacheFlush.cpp | 37 +++++----- mlir/tools/rocmlir-tuning-driver/CacheFlush.h | 8 ++- .../rocmlir-tuning-driver.cpp | 34 +++++---- .../unittests/Dialect/Rock/AmdArchDbTests.cpp | 22 ++++++ mlir/utils/performance/perfRunner.py | 69 ++++++++++++------- mlir/utils/performance/tuningRunner.py | 22 ++++-- 8 files changed, 176 insertions(+), 61 deletions(-) diff --git a/mlir/include/mlir/Dialect/Rock/IR/AmdArchDb.h b/mlir/include/mlir/Dialect/Rock/IR/AmdArchDb.h index e192d553a39c..56d4403bea19 100644 --- a/mlir/include/mlir/Dialect/Rock/IR/AmdArchDb.h +++ b/mlir/include/mlir/Dialect/Rock/IR/AmdArchDb.h @@ -119,6 +119,10 @@ AmdArchInfo lookupArchInfo(StringRef arch); bool isDirectToLDSSupported(GemmFeatures features); bool isGlobalPrefetchSupported(StringRef arch); bool isAsyncDirectToLDSSupported(StringRef arch); + +/// Get the size in bytes of the last-level cache for this architecture (the +/// AMD Infinity Cache where present, otherwise the L2). +int64_t getLastLevelCacheSize(StringRef arch); } // namespace rock } // namespace mlir diff --git a/mlir/lib/Dialect/Rock/IR/AmdArchDb.cpp b/mlir/lib/Dialect/Rock/IR/AmdArchDb.cpp index 71c4b568215e..681b4581b898 100644 --- a/mlir/lib/Dialect/Rock/IR/AmdArchDb.cpp +++ b/mlir/lib/Dialect/Rock/IR/AmdArchDb.cpp @@ -414,6 +414,47 @@ AmdArchInfo mlir::rock::lookupArchInfo(StringRef arch) { llvm_unreachable(msg.c_str()); } +int64_t mlir::rock::getLastLevelCacheSize(StringRef arch) { + constexpr int64_t kMiB = 1024 * 1024; + + // We cannot rely on hipDeviceProp_t::l2CacheSize for last-level sizing: it + // reports the small per-XCD L2 (~4 MiB on CDNA3/CDNA4), not the last-level + // AMD Infinity Cache that actually needs to be evicted between timed runs. + // Classify by chip the same way lookupArchInfo does, but distinguish the + // generations whose Infinity Cache differs (e.g. gfx101x vs gfx103x). + auto [chip, deviceId] = parseArchString(arch); + (void)deviceId; + StringRef minor = chip.take_back(2); + StringRef major = chip.slice(0, chip.size() - 2); + + if (major == "gfx9") { + // CDNA3 (gfx942) / CDNA4 (gfx950) carry a large last-level Infinity Cache. + if (minor == "42" || minor == "50") + return 256 * kMiB; + // CDNA1 (gfx908) / CDNA2 (gfx90a) top out at a per-GCD L2. + if (minor == "08" || minor == "0a") + return 8 * kMiB; + // GCN5 / gfx906: L2 is the last level. + return 4 * kMiB; + } + if (major == "gfx10") { + // gfx103x (RDNA2) introduced the Infinity Cache; gfx101x (RDNA1) did not. + if (minor.starts_with("3")) + return 128 * kMiB; + return 4 * kMiB; + } + if (major == "gfx11") // RDNA3 + return 96 * kMiB; + if (major == "gfx12") { + // gfx1250 assumed Infinity-Cache-class; TODO confirm once AMD publishes it. + if (minor == "50") + return 256 * kMiB; + return 64 * kMiB; // RDNA4 + } + // Unknown arch: assume an Infinity-Cache-class last-level cache. + return 256 * kMiB; +} + GemmFeatures mlir::rock::AmdArchInfo::getDefaultFeatures(Type dataType) { GemmFeatures theseFeatures = defaultFeatures; bool isWmma = bitEnumContainsAll(theseFeatures, GemmFeatures::wmma); diff --git a/mlir/tools/rocmlir-tuning-driver/CacheFlush.cpp b/mlir/tools/rocmlir-tuning-driver/CacheFlush.cpp index fd82c080ab5d..10698117c3ca 100644 --- a/mlir/tools/rocmlir-tuning-driver/CacheFlush.cpp +++ b/mlir/tools/rocmlir-tuning-driver/CacheFlush.cpp @@ -172,12 +172,10 @@ class CacheFlushState { } } - LogicalResult flushL2Cache(hipStream_t stream) { + LogicalResult flushCache(hipStream_t stream, bool useLastLevelCacheSize) { std::lock_guard lock(stateMutex); - if (failed(allocL2CacheFlushBuffer())) + if (failed(allocCacheFlushBuffer(useLastLevelCacheSize))) return failure(); - if (skipL2Flush) - return success(); CHECK_HIP(hipMemsetAsync(flushBuffer.get(), 0, flushSize, stream)); return success(); } @@ -205,7 +203,6 @@ class CacheFlushState { result = failure(); } flushSize = 0; - skipL2Flush = false; #if defined(__HIP_PLATFORM_AMD__) if (failed(icacheKernel.cleanup())) result = failure(); @@ -216,17 +213,24 @@ class CacheFlushState { } private: - LogicalResult allocL2CacheFlushBuffer() { - if (flushBuffer || skipL2Flush) - return success(); - size_t l2Size = deviceProps.l2CacheSize; - if (l2Size == 0) { - llvm::errs() << "Device '" << deviceProps.name - << "' reported zero-sized L2 cache; skipping L2 flush.\n"; - skipL2Flush = true; + LogicalResult allocCacheFlushBuffer(bool useLastLevelCacheSize) { + if (flushBuffer) return success(); + if (useLastLevelCacheSize) { + // Size the flush buffer to the architecture's last-level cache. We cannot + // use hipDeviceProp_t::l2CacheSize because it only reports the small + // per-XCD L2 (~4 MiB on CDNA3/CDNA4), not the last-level AMD Infinity + // Cache that actually needs to be evicted between timed runs (256 MiB on + // MI300X/MI325X/MI350X). rock::getLastLevelCacheSize returns the right + // last-level size per arch (Infinity Cache where present, else L2). + flushSize = static_cast( + rock::getLastLevelCacheSize(deviceProps.gcnArchName)); + } else { + // Default: size the flush buffer to the L2 cache reported by the HIP + // runtime, plus a 20% margin. + size_t l2Size = static_cast(deviceProps.l2CacheSize); + flushSize = l2Size + (l2Size / 5); // 20% margin } - flushSize = l2Size + (l2Size / 5); // 20% margin void *rawBuffer = nullptr; CHECK_HIP(hipMalloc(&rawBuffer, flushSize)); flushBuffer.reset(rawBuffer); @@ -265,7 +269,6 @@ class CacheFlushState { hipDeviceProp_t deviceProps = {}; size_t flushSize = 0; HipDeviceBuffer flushBuffer; - bool skipL2Flush = false; #if defined(__HIP_PLATFORM_AMD__) static constexpr int32_t kDefaultWaveSize = 64; // https://github.com/ROCm/composable_kernel/blob/develop/include/ck_tile/host/flush_icache.hpp @@ -308,8 +311,8 @@ CacheFlushState &getState() { } // namespace -LogicalResult flushL2Cache(hipStream_t stream) { - return getState().flushL2Cache(stream); +LogicalResult flushCache(hipStream_t stream, bool useLastLevelCacheSize) { + return getState().flushCache(stream, useLastLevelCacheSize); } LogicalResult flushInstructionCache(hipStream_t stream) { diff --git a/mlir/tools/rocmlir-tuning-driver/CacheFlush.h b/mlir/tools/rocmlir-tuning-driver/CacheFlush.h index 6edb9defb7e1..dbd0f57e381e 100644 --- a/mlir/tools/rocmlir-tuning-driver/CacheFlush.h +++ b/mlir/tools/rocmlir-tuning-driver/CacheFlush.h @@ -15,10 +15,14 @@ namespace rocmlir::tuningdriver { -/// \brief Flushes the L2 cache by performing a memory write operation. +/// \brief Flushes the cache by performing a memory write operation. /// \param stream The HIP stream to use for the flush operation. +/// \param useLastLevelCacheSize When true, size the flush buffer to the +/// architecture's last-level cache (e.g. AMD Infinity Cache) instead of the +/// per-XCD L2 cache size reported by the HIP runtime. /// \return success() if the flush succeeds, failure() otherwise. -mlir::LogicalResult flushL2Cache(hipStream_t stream); +mlir::LogicalResult flushCache(hipStream_t stream, + bool useLastLevelCacheSize = false); /// \brief Flushes the instruction cache to ensure that any modified code is /// visible to the device. diff --git a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp index 50cff03b2a31..60bf0b36aff1 100644 --- a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp +++ b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp @@ -200,6 +200,14 @@ static llvm::cl::opt gpuRunTimeout( "code."), llvm::cl::value_desc("seconds"), llvm::cl::init(0)); +static llvm::cl::opt flushLastLevelCache( + "flush-last-level-cache", + llvm::cl::desc( + "Size the cache-flush buffer to the architecture's last-level cache " + "(e.g. AMD Infinity Cache) instead of the per-XCD L2 cache size " + "reported by the HIP runtime. Defaults to the L2 cache size."), + llvm::cl::init(false)); + // Ripped out of JitRunner.cpp static OwningOpRef parseMLIRInput(StringRef inputFilename, MLIRContext *context) { @@ -341,6 +349,7 @@ struct BenchmarkParams { std::string benchmarkConfig; bool waitForCompiles; unsigned gpuRunTimeoutSec; + bool flushLastLevelCache; }; enum class CompilationStatus { @@ -410,14 +419,13 @@ struct ThreadResources { bool isValid() const { return sourceModule && *sourceModule; } }; -static LogicalResult -measureKernel(unsigned iterations, hipStream_t stream, - const std::vector &functions, - ArrayRef blockSizes, ArrayRef gridSizes, - std::vector &argPointers, - std::vector &measurements, - const std::optional &gpuRunDeadline, - unsigned timeoutSec, StringRef perfConfig) { +static LogicalResult measureKernel( + unsigned iterations, hipStream_t stream, + const std::vector &functions, ArrayRef blockSizes, + ArrayRef gridSizes, std::vector &argPointers, + std::vector &measurements, + const std::optional &gpuRunDeadline, unsigned timeoutSec, + StringRef perfConfig, bool useLastLevelCacheSize) { // Pre-allocate one event pair per iteration so we can record them all in a // tight loop and synchronize only once at the end. This matches Triton's // do_bench, which minimizes host-side overhead between launches (no @@ -446,7 +454,7 @@ measureKernel(unsigned iterations, hipStream_t stream, if (failed(flushInstructionCache(stream))) { return failure(); } - if (failed(flushL2Cache(stream))) { + if (failed(flushCache(stream, useLastLevelCacheSize))) { return failure(); } @@ -545,7 +553,7 @@ benchmarkKernels(ArrayRef binaries, // includes the cache clear. if (failed(flushInstructionCache(stream))) return failure(); - if (failed(flushL2Cache(stream))) + if (failed(flushCache(stream, params.flushLastLevelCache))) return failure(); for (auto [func, blockSize, gridSize] : llvm::zip(functions, blockSizes, gridSizes)) { @@ -599,7 +607,8 @@ benchmarkKernels(ArrayRef binaries, if (failed(measureKernel(iterations, stream, functions, blockSizes, gridSizes, argPointers, measurements, gpuRunDeadline, - params.gpuRunTimeoutSec, perfConfig))) + params.gpuRunTimeoutSec, perfConfig, + params.flushLastLevelCache))) return failure(); if (params.showAllMeasurements) { @@ -782,7 +791,8 @@ static LogicalResult runTuningLoop(ModuleOp source) { numCompileThreads, benchmarkConfig, waitForCompiles, - gpuRunTimeout}; + gpuRunTimeout, + flushLastLevelCache}; rock::TuningParamSetKind effectiveKind = benchmarkParams.tuningSpaceKind; unsigned numTuningIterations = rock::getNumberOfIterations(effectiveKind); diff --git a/mlir/unittests/Dialect/Rock/AmdArchDbTests.cpp b/mlir/unittests/Dialect/Rock/AmdArchDbTests.cpp index f0f04b0c32ed..95f84fbd9013 100644 --- a/mlir/unittests/Dialect/Rock/AmdArchDbTests.cpp +++ b/mlir/unittests/Dialect/Rock/AmdArchDbTests.cpp @@ -65,3 +65,25 @@ TEST_P(NativeArchTest, NativeArchInfoMatchesPresetInfo) { INSTANTIATE_TEST_SUITE_P(NativeArchTests, NativeArchTest, NativeArchTest::getDeviceIds()); + +// --- getLastLevelCacheSize --- + +TEST(AmdArchDbTest, LastLevelCacheSize) { + constexpr int64_t kMiB = 1024 * 1024; + EXPECT_EQ(getLastLevelCacheSize("gfx906"), 4 * kMiB); + EXPECT_EQ(getLastLevelCacheSize("gfx908"), 8 * kMiB); + EXPECT_EQ(getLastLevelCacheSize("gfx90a"), 8 * kMiB); + EXPECT_EQ(getLastLevelCacheSize("gfx942"), 256 * kMiB); + EXPECT_EQ(getLastLevelCacheSize("gfx950"), 256 * kMiB); + EXPECT_EQ(getLastLevelCacheSize("gfx1010"), 4 * kMiB); + EXPECT_EQ(getLastLevelCacheSize("gfx1030"), 128 * kMiB); + EXPECT_EQ(getLastLevelCacheSize("gfx1100"), 96 * kMiB); + EXPECT_EQ(getLastLevelCacheSize("gfx1200"), 64 * kMiB); + EXPECT_EQ(getLastLevelCacheSize("gfx1250"), 256 * kMiB); +} + +TEST(AmdArchDbTest, LastLevelCacheSizeWithTriple) { + constexpr int64_t kMiB = 1024 * 1024; + EXPECT_EQ(getLastLevelCacheSize("amdgcn-amd-amdhsa:gfx942"), 256 * kMiB); + EXPECT_EQ(getLastLevelCacheSize("amdgcn-amd-amdhsa:gfx906:xnack-"), 4 * kMiB); +} diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py index dc602cfde6f2..12e2d5ba8027 100644 --- a/mlir/utils/performance/perfRunner.py +++ b/mlir/utils/performance/perfRunner.py @@ -62,11 +62,12 @@ # Time budgets (ms) for the tuning-driver benchmark. The number of warmup and # measured iterations is derived from these budgets and the estimated per-launch -# runtime (Triton do_bench style). Deliberately stricter than Triton's defaults -# (25/100) to get more stable performance numbers. -BENCH_WARMUP_MS = 50 -BENCH_REP_MS = 200 -SLEEP_US = 1000 # 1 ms +# runtime (Triton do_bench style). These mirror Triton's do_bench defaults. +# tuningRunner imports these so that inference (benchmark) numbers stay consistent +# with tuning numbers for the same perfConfig. +TUNE_WARMUP_MS = 25 +TUNE_REP_MS = 100 +SLEEP_US = 100 # 0.1 ms FILTER_LAYOUT_MAP = {'N': 'k', 'C': 'c', 'H': 'y', 'W': 'x', 'G': 'g', '0': '0', '1': '1'} INPUT_LAYOUT_MAP = {'N': 'n', 'C': 'c', 'H': 'h', 'W': 'w', 'G': 'g', '0': '0', '1': '1'} @@ -2010,6 +2011,7 @@ def run_config_with_mlir(config: PerfConfiguration, arch, rocmlir_gen_flags, use_rocprof=False, + flush_last_level_cache=False, debug=True): # remove the result file generated by rocprof in previous benchmarking if os.path.exists(get_profiler_output_path(arch, BENCHMARKING_STATS_FILE_NAME)): @@ -2035,9 +2037,12 @@ def run_config_with_mlir(config: PerfConfiguration, print("Using HIP timing for benchmarking") tuning_driver_command = [ paths.mlir_paths.rocmlir_tuning_driver_path, f'--benchmark-config={config.perfconfig}', - f'--rep={BENCH_REP_MS}', f'--warmup={BENCH_WARMUP_MS}', f'--sleep-us={SLEEP_US}', - '--use-median', '-' + f'--rep={TUNE_REP_MS}', f'--warmup={TUNE_WARMUP_MS}', f'--sleep-us={SLEEP_US}', + '--use-median' ] + if flush_last_level_cache: + tuning_driver_command.append("--flush-last-level-cache") + tuning_driver_command.append('-') outs, noerr = run_pipeline([rocmlir_gen_cmd.split(), tuning_driver_command]) if noerr: try: @@ -2077,7 +2082,8 @@ def benchmark_mlir(commandline, num_chiplets, tuning_db: MaybeTuningDb, rocmlir_gen_flags, - use_rocprof=False): + use_rocprof=False, + flush_last_level_cache=False): config = conf_class.from_command_line(commandline, arch, num_cu, num_chiplets) config_str = config.to_command_line() if tuning_db: @@ -2086,7 +2092,8 @@ def benchmark_mlir(commandline, else: # Tuning DB present but doesn't contain config, return N/A return config.table_entry(np.nan) - nanoseconds = run_config_with_mlir(config, paths, arch, rocmlir_gen_flags, use_rocprof) + nanoseconds = run_config_with_mlir(config, paths, arch, rocmlir_gen_flags, use_rocprof, + flush_last_level_cache) return config.table_entry(nanoseconds) @@ -2100,22 +2107,26 @@ def generate_performance_results(configs, tuning_db: MaybeTuningDb, quick_tuning_db: MaybeTuningDb, rocmlir_gen_flags, - use_rocprof=False): + use_rocprof=False, + flush_last_level_cache=False): # Never pass tuning DB to this run mlir_df = pd.DataFrame( benchmark_mlir(test_vector.split(sep=' '), conf_class, paths, arch, num_cu, num_chiplets, - None, rocmlir_gen_flags, use_rocprof) for test_vector in configs) + None, rocmlir_gen_flags, use_rocprof, flush_last_level_cache) + for test_vector in configs) tuned_df = None if tuning_db: tuned_df = pd.DataFrame( - benchmark_mlir(test_vector.split(sep=' '), conf_class, paths, arch, num_cu, - num_chiplets, tuning_db, rocmlir_gen_flags, use_rocprof) + benchmark_mlir(test_vector.split( + sep=' '), conf_class, paths, arch, num_cu, num_chiplets, tuning_db, + rocmlir_gen_flags, use_rocprof, flush_last_level_cache) for test_vector in configs) quick_tuned_df = None if quick_tuning_db: quick_tuned_df = pd.DataFrame( - benchmark_mlir(test_vector.split(sep=' '), conf_class, paths, arch, num_cu, - num_chiplets, quick_tuning_db, rocmlir_gen_flags, use_rocprof) + benchmark_mlir(test_vector.split( + sep=' '), conf_class, paths, arch, num_cu, num_chiplets, quick_tuning_db, + rocmlir_gen_flags, use_rocprof, flush_last_level_cache) for test_vector in configs) external_df = pd.DataFrame( @@ -2358,7 +2369,8 @@ def benchmark_fusion_kernels(test_dir, num_cu, num_chiplets, tuning_db: MaybeTuningDb, - use_rocprof=False): + use_rocprof=False, + flush_last_level_cache=False): all_tests = [] # filename, test_vector, fut_name perf_results = {} # associate test_vector to config and performances chip = GFX_CHIP_RE.search(arch).group(0) @@ -2433,7 +2445,8 @@ def benchmark_fusion_kernels(test_dir, continue # Run gemm or conv op with the same configuration - nanoseconds = run_config_with_mlir(config, paths, arch, '', use_rocprof) + nanoseconds = run_config_with_mlir(config, paths, arch, '', use_rocprof, + flush_last_level_cache) one_entry['MLIR TFlops'] = config.compute_tflops(nanoseconds) one_entry['Fusion/MLIR'] = one_entry['TFlops'] / one_entry['MLIR TFlops'] one_entry['FileName'] = filename @@ -2663,6 +2676,14 @@ def main(args=None): action="store_true", help="Use rocprof instead of rocmlir-tuning-driver to collect performance data") + parser.add_argument( + "--flush-last-level-cache", + action='store_true', + default=False, + help= + "Size the cache-flush buffer to the architecture's last-level cache (e.g. AMD Infinity Cache) instead of the per-XCD L2 cache size reported by the HIP runtime. Defaults to the L2 cache size." + ) + parsed_args = parser.parse_args(args) rocmlir_gen_flags = '' @@ -2736,7 +2757,7 @@ def main(args=None): # batch benchmark with MLIR and MIOpen. generate_performance_results(configs, conf_class, paths, arch, num_cu, num_chiplets, tuning_db, quick_tuning_db, rocmlir_gen_flags, - parsed_args.use_rocprof) + parsed_args.use_rocprof, parsed_args.flush_last_level_cache) elif parsed_args.tuning: tune_mlir_kernels(configs, arch, num_cu, num_chiplets) elif optype == Operation.FUSION: @@ -2744,13 +2765,14 @@ def main(args=None): raise RuntimeError("MLIR build dir was not provided/found") else: benchmark_fusion_kernels(parsed_args.test_dir, paths, arch, num_cu, num_chiplets, - tuning_db, parsed_args.use_rocprof) + tuning_db, parsed_args.use_rocprof, + parsed_args.flush_last_level_cache) else: if parsed_args.batch_mlir: df = pd.DataFrame( benchmark_mlir(test_vector.split(sep=' '), conf_class, paths, arch, num_cu, - num_chiplets, tuning_db, rocmlir_gen_flags, parsed_args.use_rocprof) - for test_vector in configs) + num_chiplets, tuning_db, rocmlir_gen_flags, parsed_args.use_rocprof, + parsed_args.flush_last_level_cache) for test_vector in configs) elif parsed_args.batch_external: df = pd.DataFrame( conf_class.benchmark_external(test_vector.split( @@ -2769,13 +2791,14 @@ def main(args=None): df = pd.DataFrame([ benchmark_mlir(parsed_args.config, conf_class, paths, arch, num_cu, num_chiplets, tuning_db, rocmlir_gen_flags, - parsed_args.use_rocprof) + parsed_args.use_rocprof, parsed_args.flush_last_level_cache) ]) else: df = pd.DataFrame([ benchmark_mlir(config.split(), conf_class, paths, arch, num_cu, num_chiplets, tuning_db, rocmlir_gen_flags, - parsed_args.use_rocprof) for config in configs + parsed_args.use_rocprof, parsed_args.flush_last_level_cache) + for config in configs ]) df.to_csv(parsed_args.filename) with pd.option_context('display.precision', reportUtils.ROUND_DIGITS): diff --git a/mlir/utils/performance/tuningRunner.py b/mlir/utils/performance/tuningRunner.py index 6326dfc48429..0df02a046255 100755 --- a/mlir/utils/performance/tuningRunner.py +++ b/mlir/utils/performance/tuningRunner.py @@ -59,6 +59,9 @@ GemmGemmConfiguration, Paths, PerfConfiguration, + SLEEP_US, + TUNE_REP_MS, + TUNE_WARMUP_MS, canonicalize_config, ) @@ -69,13 +72,6 @@ # rocmlir-gen host-harness kernel repeat count (--kernel-repeats, used with -ph). MLIR_N_REPEATS = 10 -# Time budgets (ms) for the tuning-driver benchmark. The number of warmup and -# measured iterations is derived from these budgets and the estimated per-launch -# runtime (Triton do_bench style). These mirror Triton's do_bench defaults. -TUNE_WARMUP_MS = 25 -TUNE_REP_MS = 100 -SLEEP_US = 100 # 0.1 ms - # A GPU run timeout is different from the outer tuning subprocess timeout: an # in-process kernel may have hung and left the HIP context untrustworthy, so the # tuning driver exits the whole process with this distinct code. @@ -202,6 +198,7 @@ class Options: gpu_ids: List[int] num_cpus: Optional[int] wait_for_compiles: bool + flush_last_level_cache: bool timeout: Optional[int] verify_timeout: Optional[int] gpu_run_timeout: int @@ -1465,6 +1462,8 @@ def tune_config(test_vector: str, conf_class: type, paths: Paths, options: Optio f"--wait-for-compiles={options.wait_for_compiles}", f"--gpu-run-timeout={options.gpu_run_timeout}", ] + if options.flush_last_level_cache: + tuning_driver_args.append("--flush-last-level-cache") env = make_isolated_gpu_env(gpu_id) @@ -2140,6 +2139,14 @@ def parse_arguments(gpu_topology: GpuTopology, "Wait for all compilation tasks to complete before starting tuning. Useful for systems with shared CPU/GPU memory (e.g., APUs)." ) + parser.add_argument( + "--flush-last-level-cache", + action='store_true', + default=False, + help= + "Size the cache-flush buffer to the architecture's last-level cache (e.g. AMD Infinity Cache) instead of the per-XCD L2 cache size reported by the HIP runtime. Defaults to the L2 cache size." + ) + parser.add_argument("-s", "--status", action='store_true', @@ -2242,6 +2249,7 @@ def main(args=None): gpu_ids=parsed_args.gpus, num_cpus=parsed_args.num_cpus, wait_for_compiles=parsed_args.wait_for_compiles, + flush_last_level_cache=parsed_args.flush_last_level_cache, timeout=parsed_args.timeout, verify_timeout=parsed_args.verify_timeout, gpu_run_timeout=parsed_args.gpu_run_timeout) From ca442e442375bb50da9bf91727f6e51bd394ca6f Mon Sep 17 00:00:00 2001 From: Daniel Hernandez Date: Tue, 18 Aug 2026 11:01:30 +0200 Subject: [PATCH 3/6] Addressing PR comments --- .../rocmlir-tuning-driver/rocmlir-tuning-driver.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp index 60bf0b36aff1..5bc9c2f7886d 100644 --- a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp +++ b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp @@ -432,7 +432,7 @@ static LogicalResult measureKernel( // per-iteration synchronization). std::vector startEvents(iterations, nullptr); std::vector stopEvents(iterations, nullptr); - auto eventCleanup = llvm::make_scope_exit([&]() { + llvm::scope_exit eventCleanup([&]() { for (hipEvent_t event : startEvents) { if (event) (void)hipEventDestroy(event); @@ -499,7 +499,7 @@ benchmarkKernels(ArrayRef binaries, // Load all modules once to reduce overhead std::vector modules; std::vector functions; - auto moduleCleanup = llvm::make_scope_exit([&]() { + llvm::scope_exit moduleCleanup([&]() { for (hipModule_t mod : modules) { if (!mod) continue; @@ -522,7 +522,7 @@ benchmarkKernels(ArrayRef binaries, } // Sleep guard to avoid GPU throttling - auto sleepGuard = llvm::make_scope_exit([¶ms] { + llvm::scope_exit sleepGuard([¶ms] { if (params.sleepUs > 0) { std::this_thread::sleep_for(std::chrono::microseconds(params.sleepUs)); } @@ -737,7 +737,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: " @@ -746,7 +746,7 @@ static LogicalResult runTuningLoop(ModuleOp source) { }); std::vector 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) @@ -1015,7 +1015,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) { From 27410379e5a566688af8b13c719948ef97120f21 Mon Sep 17 00:00:00 2001 From: Daniel Hernandez Date: Tue, 18 Aug 2026 11:11:09 +0200 Subject: [PATCH 4/6] add --legacy-benchmark-mode to run the legacy benchmark mode --- .../rocmlir-tuning-driver.cpp | 359 ++++++++++++++---- mlir/utils/performance/perfRunner.py | 62 ++- mlir/utils/performance/tuningRunner.py | 12 + 3 files changed, 344 insertions(+), 89 deletions(-) diff --git a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp index 5bc9c2f7886d..c0e028b4f09a 100644 --- a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp +++ b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp @@ -147,6 +147,29 @@ static llvm::cl::opt warmup( "per-launch runtime (Triton do_bench style)."), llvm::cl::value_desc("warmup milliseconds"), llvm::cl::init(25)); +static llvm::cl::opt legacyBenchmarkMode( + "legacy-benchmark-mode", + llvm::cl::desc( + "Use the legacy rocMLIR benchmarking method (fixed iteration counts " + "with a small-vs-large-kernel split) instead of the default " + "Triton do_bench-style time-budget measurement. Kept for " + "apples-to-apples comparison against older rocMLIR versions. In this " + "mode --num-iterations/--warmup-iterations control the run counts and " + "--rep/--warmup are ignored."), + llvm::cl::init(false)); + +static llvm::cl::opt numIterations( + "num-iterations", + llvm::cl::desc("Number of times to run each kernel for averaging (only " + "used with --legacy-benchmark-mode)"), + llvm::cl::value_desc("number of runs"), llvm::cl::init(100)); + +static llvm::cl::opt warmupIterations( + "warmup-iterations", + llvm::cl::desc( + "Number of warmup runs (only used with --legacy-benchmark-mode)"), + llvm::cl::value_desc("number of warmup runs"), llvm::cl::init(10)); + static llvm::cl::opt useMedian("use-median", llvm::cl::desc("Use median of runs instead of mean for timing " @@ -350,6 +373,12 @@ struct BenchmarkParams { bool waitForCompiles; unsigned gpuRunTimeoutSec; bool flushLastLevelCache; + // Legacy benchmarking path (fixed iteration counts + small/large kernel + // split). When false, the default Triton do_bench-style measurement is used + // and numIterations/warmupIterations are ignored. + bool legacyBenchmarkMode; + unsigned numIterations; + unsigned warmupIterations; }; enum class CompilationStatus { @@ -419,6 +448,98 @@ struct ThreadResources { bool isValid() const { return sourceModule && *sourceModule; } }; +// Legacy small-kernel path: run all iterations back-to-back and time the whole +// batch with a single CPU timer. Only used with --legacy-benchmark-mode. +static LogicalResult +measureSmallKernel(unsigned iterations, hipStream_t stream, + const std::vector &functions, + ArrayRef blockSizes, ArrayRef gridSizes, + std::vector &argPointers, + std::vector &measurements, double &smallKernelCpuMs, + bool benchmarkMode, bool useLastLevelCacheSize, + const std::optional &gpuRunDeadline, + unsigned timeoutSec, StringRef perfConfig) { + // Special case for small kernels, where we measure the time for all kernels + // at once, using CPU timers. + auto iterationStart = std::chrono::steady_clock::now(); + for (unsigned iter = 0; iter < iterations; ++iter) { + // Do not flush caches in benchmark mode, as we do not want to + // time the cache flush (it's okay if we are running in tuning mode). + if (!benchmarkMode) { + if (failed(flushInstructionCache(stream))) { + return failure(); + } + if (failed(flushCache(stream, useLastLevelCacheSize))) { + return failure(); + } + } + for (auto [func, blockSize, gridSize] : + llvm::zip(functions, blockSizes, gridSizes)) { + HIPCHECK(hipExtModuleLaunchKernel( + func, gridSize * blockSize, 1, 1, blockSize, 1, 1, 0, stream, + argPointers.data(), nullptr, nullptr, nullptr)); + } + } + + if (failed(synchronizeStreamWithTimeout(stream, gpuRunDeadline, timeoutSec, + perfConfig, "measurement"))) + return failure(); + smallKernelCpuMs = std::chrono::duration( + std::chrono::steady_clock::now() - iterationStart) + .count(); + measurements.push_back(smallKernelCpuMs / iterations); + return success(); +} + +// Legacy large-kernel path: time each iteration individually with GPU events, +// synchronizing after every launch. Only used with --legacy-benchmark-mode. +static LogicalResult measureLargeKernel( + unsigned iterations, hipStream_t stream, + const std::vector &functions, ArrayRef blockSizes, + ArrayRef gridSizes, std::vector &argPointers, + std::vector &measurements, bool useLastLevelCacheSize, + const std::optional &gpuRunDeadline, unsigned timeoutSec, + StringRef perfConfig) { + // Measure runs normally. + for (unsigned iter = 0; iter < iterations; ++iter) { + if (failed(flushInstructionCache(stream))) { + return failure(); + } + if (failed(flushCache(stream, useLastLevelCacheSize))) { + return failure(); + } + + double totalMilliseconds = 0.0; + + for (auto [func, blockSize, gridSize] : + llvm::zip(functions, blockSizes, gridSizes)) { + hipEvent_t startEvent, stopEvent; + HIPCHECK(hipEventCreate(&startEvent)); + HIPCHECK(hipEventCreate(&stopEvent)); + + HIPCHECK(hipExtModuleLaunchKernel( + func, gridSize * blockSize, 1, 1, blockSize, 1, 1, 0, stream, + argPointers.data(), nullptr, startEvent, stopEvent)); + if (failed(synchronizeStreamWithTimeout( + stream, gpuRunDeadline, timeoutSec, perfConfig, "measurement"))) + return failure(); + + float currentMilliseconds = 0.0; + HIPCHECK( + hipEventElapsedTime(¤tMilliseconds, startEvent, stopEvent)); + + HIPCHECK(hipEventDestroy(stopEvent)); + HIPCHECK(hipEventDestroy(startEvent)); + + totalMilliseconds += static_cast(currentMilliseconds); + } + + measurements.push_back(totalMilliseconds); + } + + return success(); +} + static LogicalResult measureKernel( unsigned iterations, hipStream_t stream, const std::vector &functions, ArrayRef blockSizes, @@ -534,27 +655,151 @@ benchmarkKernels(ArrayRef binaries, std::optional gpuRunDeadline = makeTimeoutDeadline(params.gpuRunTimeoutSec); - // Estimate the per-launch runtime so we can size warmup/benchmark iteration - // counts from the requested time budgets (Triton do_bench style). We time a - // handful of launches (flushing caches between them) using a single event - // pair. - constexpr unsigned estimateRuns = 5; - double estimateMs = 0.0; - { - hipEvent_t startEvent, stopEvent; - HIPCHECK(hipEventCreate(&startEvent)); - HIPCHECK(hipEventCreate(&stopEvent)); - - HIPCHECK(hipEventRecord(startEvent, stream)); - for (unsigned iter = 0; iter < estimateRuns; ++iter) { - // The cache flushes are inside the timed window here (unlike the actual - // measurement loop, which flushes before recording the start event). This - // is intentional, to match Triton's do_bench, whose estimate loop also - // includes the cache clear. - if (failed(flushInstructionCache(stream))) + std::vector measurements; + double smallKernelCpuMs = 0.0; + bool isSmallKernel = false; + unsigned smallKernelIters = 0; + + if (params.legacyBenchmarkMode) { + // Legacy benchmarking: fixed --num-iterations/--warmup-iterations counts + // with a small-vs-large-kernel split. Kept for apples-to-apples comparison + // with older rocMLIR versions. + bool benchmarkMode = !params.benchmarkConfig.empty(); + unsigned iterations = params.numIterations; + + if (params.warmupIterations > 0) { + // Warmup run. We measure the warmup to get an estimate of the kernel + // runtime. We will use this estimate to determine if the kernel is small + // or not. + double totalMillisecondsWarmup = 0.0; + for (unsigned iter = 0; iter < params.warmupIterations; ++iter) { + for (auto [func, blockSize, gridSize] : + llvm::zip(functions, blockSizes, gridSizes)) { + hipEvent_t startEvent, stopEvent; + HIPCHECK(hipEventCreate(&startEvent)); + HIPCHECK(hipEventCreate(&stopEvent)); + + HIPCHECK(hipExtModuleLaunchKernel( + func, gridSize * blockSize, 1, 1, blockSize, 1, 1, 0, stream, + argPointers.data(), nullptr, startEvent, stopEvent)); + + if (failed(synchronizeStreamWithTimeout(stream, gpuRunDeadline, + params.gpuRunTimeoutSec, + perfConfig, "warmup"))) + return failure(); + + float currentMilliseconds = 0.0; + HIPCHECK( + hipEventElapsedTime(¤tMilliseconds, startEvent, stopEvent)); + + HIPCHECK(hipEventDestroy(stopEvent)); + HIPCHECK(hipEventDestroy(startEvent)); + + // hipEventElapsedTime seemingly can return negative values for fast + // kernels due to GPU clock precision issues. This is extremely + // relevant when we have a small number of warmup iterations (e.g., 1) + // for small kernels. Clamp to the documented resolution of ~1 + // microsecond (0.001 ms) if this is the case. + if (currentMilliseconds < 0.0f) { + constexpr float minMeasurableMs = 0.001f; + currentMilliseconds = minMeasurableMs; + } + + totalMillisecondsWarmup += static_cast(currentMilliseconds); + } + } + totalMillisecondsWarmup /= params.warmupIterations; + assert(totalMillisecondsWarmup >= 0.0f && + "totalMillisecondsWarmup must be greater than 0"); + + // We want to get at least 1ms of kernel execution time + // (counting all iterations), so increase the number of iterations + // if necessary. + constexpr float minTotalMilliseconds = 1.0f; + iterations = std::max( + iterations, static_cast(std::ceil( + minTotalMilliseconds / totalMillisecondsWarmup))); + + // Depending on the runtime of the kernel, we will use a different + // approach to measure the runs. We consider a kernel to be small if a + // single iteration takes less than 1ms to run. + constexpr float smallKernelThreshold = 1.0f; + isSmallKernel = totalMillisecondsWarmup < smallKernelThreshold; + } + + smallKernelIters = iterations; + if (isSmallKernel) { + if (failed(measureSmallKernel(iterations, stream, functions, blockSizes, + gridSizes, argPointers, measurements, + smallKernelCpuMs, benchmarkMode, + params.flushLastLevelCache, gpuRunDeadline, + params.gpuRunTimeoutSec, perfConfig))) return failure(); - if (failed(flushCache(stream, params.flushLastLevelCache))) + } else { + if (failed(measureLargeKernel(iterations, stream, functions, blockSizes, + gridSizes, argPointers, measurements, + params.flushLastLevelCache, gpuRunDeadline, + params.gpuRunTimeoutSec, perfConfig))) + return failure(); + } + } else { + // Estimate the per-launch runtime so we can size warmup/benchmark iteration + // counts from the requested time budgets (Triton do_bench style). We time a + // handful of launches (flushing caches between them) using a single event + // pair. + constexpr unsigned estimateRuns = 5; + double estimateMs = 0.0; + { + hipEvent_t startEvent, stopEvent; + HIPCHECK(hipEventCreate(&startEvent)); + HIPCHECK(hipEventCreate(&stopEvent)); + + HIPCHECK(hipEventRecord(startEvent, stream)); + for (unsigned iter = 0; iter < estimateRuns; ++iter) { + // The cache flushes are inside the timed window here (unlike the actual + // measurement loop, which flushes before recording the start event). + // This is intentional, to match Triton's do_bench, whose estimate loop + // also includes the cache clear. + if (failed(flushInstructionCache(stream))) + return failure(); + if (failed(flushCache(stream, params.flushLastLevelCache))) + return failure(); + for (auto [func, blockSize, gridSize] : + llvm::zip(functions, blockSizes, gridSizes)) { + HIPCHECK(hipExtModuleLaunchKernel( + func, gridSize * blockSize, 1, 1, blockSize, 1, 1, 0, stream, + argPointers.data(), nullptr, nullptr, nullptr)); + } + } + HIPCHECK(hipEventRecord(stopEvent, stream)); + if (failed(synchronizeStreamWithTimeout(stream, gpuRunDeadline, + params.gpuRunTimeoutSec, + perfConfig, "warmup"))) return failure(); + + float elapsedMs = 0.0; + HIPCHECK(hipEventElapsedTime(&elapsedMs, startEvent, stopEvent)); + HIPCHECK(hipEventDestroy(stopEvent)); + HIPCHECK(hipEventDestroy(startEvent)); + + estimateMs = static_cast(elapsedMs) / estimateRuns; + // hipEventElapsedTime can return tiny/negative values for very fast + // kernels due to GPU clock precision. Clamp to the documented ~1 + // microsecond resolution (0.001 ms) to avoid divide-by-zero / overflow + // below. + constexpr double minMeasurableMs = 0.001; + if (estimateMs < minMeasurableMs) + estimateMs = minMeasurableMs; + } + + // Derive iteration counts from the time budgets, like Triton's do_bench. + unsigned nWarmup = std::max( + 1, static_cast(params.warmupMs / estimateMs)); + unsigned iterations = + std::max(1, static_cast(params.repMs / estimateMs)); + + // Warm-up (untimed): just run the kernel chain nWarmup times. + for (unsigned iter = 0; iter < nWarmup; ++iter) { for (auto [func, blockSize, gridSize] : llvm::zip(functions, blockSizes, gridSizes)) { HIPCHECK(hipExtModuleLaunchKernel( @@ -562,68 +807,43 @@ benchmarkKernels(ArrayRef binaries, argPointers.data(), nullptr, nullptr, nullptr)); } } - HIPCHECK(hipEventRecord(stopEvent, stream)); if (failed(synchronizeStreamWithTimeout(stream, gpuRunDeadline, params.gpuRunTimeoutSec, perfConfig, "warmup"))) return failure(); - float elapsedMs = 0.0; - HIPCHECK(hipEventElapsedTime(&elapsedMs, startEvent, stopEvent)); - HIPCHECK(hipEventDestroy(stopEvent)); - HIPCHECK(hipEventDestroy(startEvent)); - - estimateMs = static_cast(elapsedMs) / estimateRuns; - // hipEventElapsedTime can return tiny/negative values for very fast kernels - // due to GPU clock precision. Clamp to the documented ~1 microsecond - // resolution (0.001 ms) to avoid divide-by-zero / overflow below. - constexpr double minMeasurableMs = 0.001; - if (estimateMs < minMeasurableMs) - estimateMs = minMeasurableMs; - } - - // Derive iteration counts from the time budgets, like Triton's do_bench. - unsigned nWarmup = std::max( - 1, static_cast(params.warmupMs / estimateMs)); - unsigned iterations = - std::max(1, static_cast(params.repMs / estimateMs)); - - // Warm-up (untimed): just run the kernel chain nWarmup times. - for (unsigned iter = 0; iter < nWarmup; ++iter) { - for (auto [func, blockSize, gridSize] : - llvm::zip(functions, blockSizes, gridSizes)) { - HIPCHECK(hipExtModuleLaunchKernel( - func, gridSize * blockSize, 1, 1, blockSize, 1, 1, 0, stream, - argPointers.data(), nullptr, nullptr, nullptr)); - } + // Measure runs + if (failed(measureKernel(iterations, stream, functions, blockSizes, + gridSizes, argPointers, measurements, + gpuRunDeadline, params.gpuRunTimeoutSec, + perfConfig, params.flushLastLevelCache))) + return failure(); } - if (failed(synchronizeStreamWithTimeout(stream, gpuRunDeadline, - params.gpuRunTimeoutSec, perfConfig, - "warmup"))) - return failure(); - - // Measure runs - std::vector measurements; - - if (failed(measureKernel(iterations, stream, functions, blockSizes, gridSizes, - argPointers, measurements, gpuRunDeadline, - params.gpuRunTimeoutSec, perfConfig, - params.flushLastLevelCache))) - return failure(); if (params.showAllMeasurements) { - llvm::outs() << "["; - for (size_t i = 0; i < measurements.size(); ++i) { - if (i > 0) - llvm::outs() << ","; - llvm::outs() << measurements[i]; + if (isSmallKernel) { + llvm::outs() << "{\"total_cpu_time\":" << smallKernelCpuMs + << ",\"iterations\":" << smallKernelIters << "}\t"; + } else { + llvm::outs() << "["; + for (size_t i = 0; i < measurements.size(); ++i) { + if (i > 0) + llvm::outs() << ","; + llvm::outs() << measurements[i]; + } + llvm::outs() << "]\t"; } - llvm::outs() << "]\t"; } std::sort(measurements.begin(), measurements.end()); if (params.showStats) { + // The legacy small-kernel path uses a single CPU timer, so only the + // aggregate CPU time is available (no per-run min/max/median). + if (isSmallKernel) { + llvm::outs() << "{\"total_cpu_time\":" << smallKernelCpuMs + << ",\"iterations\":" << smallKernelIters << "}\t"; + } if (measurements.size() > 1) { float median = computeMedian(measurements); float min = measurements.front(); @@ -792,7 +1012,10 @@ static LogicalResult runTuningLoop(ModuleOp source) { benchmarkConfig, waitForCompiles, gpuRunTimeout, - flushLastLevelCache}; + flushLastLevelCache, + legacyBenchmarkMode, + numIterations, + warmupIterations}; rock::TuningParamSetKind effectiveKind = benchmarkParams.tuningSpaceKind; unsigned numTuningIterations = rock::getNumberOfIterations(effectiveKind); diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py index 12e2d5ba8027..cb1f1058c099 100644 --- a/mlir/utils/performance/perfRunner.py +++ b/mlir/utils/performance/perfRunner.py @@ -2012,6 +2012,7 @@ def run_config_with_mlir(config: PerfConfiguration, rocmlir_gen_flags, use_rocprof=False, flush_last_level_cache=False, + legacy_benchmark_mode=False, debug=True): # remove the result file generated by rocprof in previous benchmarking if os.path.exists(get_profiler_output_path(arch, BENCHMARKING_STATS_FILE_NAME)): @@ -2042,6 +2043,8 @@ def run_config_with_mlir(config: PerfConfiguration, ] if flush_last_level_cache: tuning_driver_command.append("--flush-last-level-cache") + if legacy_benchmark_mode: + tuning_driver_command.append("--legacy-benchmark-mode") tuning_driver_command.append('-') outs, noerr = run_pipeline([rocmlir_gen_cmd.split(), tuning_driver_command]) if noerr: @@ -2083,7 +2086,8 @@ def benchmark_mlir(commandline, tuning_db: MaybeTuningDb, rocmlir_gen_flags, use_rocprof=False, - flush_last_level_cache=False): + flush_last_level_cache=False, + legacy_benchmark_mode=False): config = conf_class.from_command_line(commandline, arch, num_cu, num_chiplets) config_str = config.to_command_line() if tuning_db: @@ -2093,7 +2097,7 @@ def benchmark_mlir(commandline, return config.table_entry(np.nan) nanoseconds = run_config_with_mlir(config, paths, arch, rocmlir_gen_flags, use_rocprof, - flush_last_level_cache) + flush_last_level_cache, legacy_benchmark_mode) return config.table_entry(nanoseconds) @@ -2108,25 +2112,27 @@ def generate_performance_results(configs, quick_tuning_db: MaybeTuningDb, rocmlir_gen_flags, use_rocprof=False, - flush_last_level_cache=False): + flush_last_level_cache=False, + legacy_benchmark_mode=False): # Never pass tuning DB to this run mlir_df = pd.DataFrame( - benchmark_mlir(test_vector.split(sep=' '), conf_class, paths, arch, num_cu, num_chiplets, - None, rocmlir_gen_flags, use_rocprof, flush_last_level_cache) + benchmark_mlir(test_vector.split( + sep=' '), conf_class, paths, arch, num_cu, num_chiplets, None, rocmlir_gen_flags, + use_rocprof, flush_last_level_cache, legacy_benchmark_mode) for test_vector in configs) tuned_df = None if tuning_db: tuned_df = pd.DataFrame( - benchmark_mlir(test_vector.split( - sep=' '), conf_class, paths, arch, num_cu, num_chiplets, tuning_db, - rocmlir_gen_flags, use_rocprof, flush_last_level_cache) + benchmark_mlir(test_vector.split(sep=' '), conf_class, paths, arch, num_cu, + num_chiplets, tuning_db, rocmlir_gen_flags, use_rocprof, + flush_last_level_cache, legacy_benchmark_mode) for test_vector in configs) quick_tuned_df = None if quick_tuning_db: quick_tuned_df = pd.DataFrame( - benchmark_mlir(test_vector.split( - sep=' '), conf_class, paths, arch, num_cu, num_chiplets, quick_tuning_db, - rocmlir_gen_flags, use_rocprof, flush_last_level_cache) + benchmark_mlir(test_vector.split(sep=' '), conf_class, paths, arch, num_cu, + num_chiplets, quick_tuning_db, rocmlir_gen_flags, use_rocprof, + flush_last_level_cache, legacy_benchmark_mode) for test_vector in configs) external_df = pd.DataFrame( @@ -2370,7 +2376,8 @@ def benchmark_fusion_kernels(test_dir, num_chiplets, tuning_db: MaybeTuningDb, use_rocprof=False, - flush_last_level_cache=False): + flush_last_level_cache=False, + legacy_benchmark_mode=False): all_tests = [] # filename, test_vector, fut_name perf_results = {} # associate test_vector to config and performances chip = GFX_CHIP_RE.search(arch).group(0) @@ -2446,7 +2453,7 @@ def benchmark_fusion_kernels(test_dir, # Run gemm or conv op with the same configuration nanoseconds = run_config_with_mlir(config, paths, arch, '', use_rocprof, - flush_last_level_cache) + flush_last_level_cache, legacy_benchmark_mode) one_entry['MLIR TFlops'] = config.compute_tflops(nanoseconds) one_entry['Fusion/MLIR'] = one_entry['TFlops'] / one_entry['MLIR TFlops'] one_entry['FileName'] = filename @@ -2684,6 +2691,14 @@ def main(args=None): "Size the cache-flush buffer to the architecture's last-level cache (e.g. AMD Infinity Cache) instead of the per-XCD L2 cache size reported by the HIP runtime. Defaults to the L2 cache size." ) + parser.add_argument( + "--legacy-benchmark-mode", + action='store_true', + default=False, + help= + "Use the legacy rocMLIR benchmarking method (fixed iteration counts with a small-vs-large-kernel split) instead of the default Triton do_bench-style time-budget measurement. Kept for apples-to-apples comparison against older rocMLIR versions." + ) + parsed_args = parser.parse_args(args) rocmlir_gen_flags = '' @@ -2757,7 +2772,8 @@ def main(args=None): # batch benchmark with MLIR and MIOpen. generate_performance_results(configs, conf_class, paths, arch, num_cu, num_chiplets, tuning_db, quick_tuning_db, rocmlir_gen_flags, - parsed_args.use_rocprof, parsed_args.flush_last_level_cache) + parsed_args.use_rocprof, parsed_args.flush_last_level_cache, + parsed_args.legacy_benchmark_mode) elif parsed_args.tuning: tune_mlir_kernels(configs, arch, num_cu, num_chiplets) elif optype == Operation.FUSION: @@ -2766,13 +2782,16 @@ def main(args=None): else: benchmark_fusion_kernels(parsed_args.test_dir, paths, arch, num_cu, num_chiplets, tuning_db, parsed_args.use_rocprof, - parsed_args.flush_last_level_cache) + parsed_args.flush_last_level_cache, + parsed_args.legacy_benchmark_mode) else: if parsed_args.batch_mlir: df = pd.DataFrame( - benchmark_mlir(test_vector.split(sep=' '), conf_class, paths, arch, num_cu, - num_chiplets, tuning_db, rocmlir_gen_flags, parsed_args.use_rocprof, - parsed_args.flush_last_level_cache) for test_vector in configs) + benchmark_mlir(test_vector.split( + sep=' '), conf_class, paths, arch, num_cu, num_chiplets, tuning_db, + rocmlir_gen_flags, parsed_args.use_rocprof, parsed_args. + flush_last_level_cache, parsed_args.legacy_benchmark_mode) + for test_vector in configs) elif parsed_args.batch_external: df = pd.DataFrame( conf_class.benchmark_external(test_vector.split( @@ -2791,14 +2810,15 @@ def main(args=None): df = pd.DataFrame([ benchmark_mlir(parsed_args.config, conf_class, paths, arch, num_cu, num_chiplets, tuning_db, rocmlir_gen_flags, - parsed_args.use_rocprof, parsed_args.flush_last_level_cache) + parsed_args.use_rocprof, parsed_args.flush_last_level_cache, + parsed_args.legacy_benchmark_mode) ]) else: df = pd.DataFrame([ benchmark_mlir(config.split(), conf_class, paths, arch, num_cu, num_chiplets, tuning_db, rocmlir_gen_flags, - parsed_args.use_rocprof, parsed_args.flush_last_level_cache) - for config in configs + parsed_args.use_rocprof, parsed_args.flush_last_level_cache, + parsed_args.legacy_benchmark_mode) for config in configs ]) df.to_csv(parsed_args.filename) with pd.option_context('display.precision', reportUtils.ROUND_DIGITS): diff --git a/mlir/utils/performance/tuningRunner.py b/mlir/utils/performance/tuningRunner.py index 0df02a046255..3bbb2494557a 100755 --- a/mlir/utils/performance/tuningRunner.py +++ b/mlir/utils/performance/tuningRunner.py @@ -199,6 +199,7 @@ class Options: num_cpus: Optional[int] wait_for_compiles: bool flush_last_level_cache: bool + legacy_benchmark_mode: bool timeout: Optional[int] verify_timeout: Optional[int] gpu_run_timeout: int @@ -1464,6 +1465,8 @@ def tune_config(test_vector: str, conf_class: type, paths: Paths, options: Optio ] if options.flush_last_level_cache: tuning_driver_args.append("--flush-last-level-cache") + if options.legacy_benchmark_mode: + tuning_driver_args.append("--legacy-benchmark-mode") env = make_isolated_gpu_env(gpu_id) @@ -2147,6 +2150,14 @@ def parse_arguments(gpu_topology: GpuTopology, "Size the cache-flush buffer to the architecture's last-level cache (e.g. AMD Infinity Cache) instead of the per-XCD L2 cache size reported by the HIP runtime. Defaults to the L2 cache size." ) + parser.add_argument( + "--legacy-benchmark-mode", + action='store_true', + default=False, + help= + "Use the legacy rocMLIR benchmarking method (fixed iteration counts with a small-vs-large-kernel split) instead of the default Triton do_bench-style time-budget measurement. Kept for apples-to-apples comparison against older rocMLIR versions." + ) + parser.add_argument("-s", "--status", action='store_true', @@ -2250,6 +2261,7 @@ def main(args=None): num_cpus=parsed_args.num_cpus, wait_for_compiles=parsed_args.wait_for_compiles, flush_last_level_cache=parsed_args.flush_last_level_cache, + legacy_benchmark_mode=parsed_args.legacy_benchmark_mode, timeout=parsed_args.timeout, verify_timeout=parsed_args.verify_timeout, gpu_run_timeout=parsed_args.gpu_run_timeout) From e659e6c3b48a89d2055162208b159ec140cef403 Mon Sep 17 00:00:00 2001 From: Daniel Hernandez Date: Tue, 18 Aug 2026 16:37:27 +0200 Subject: [PATCH 5/6] Addressing PR comments --- mlir/utils/performance/perfRunner.py | 17 ++++++++++++++++- mlir/utils/performance/tuningRunner.py | 17 ++++++++++++++++- 2 files changed, 32 insertions(+), 2 deletions(-) diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py index cb1f1058c099..2c88748fb2ab 100644 --- a/mlir/utils/performance/perfRunner.py +++ b/mlir/utils/performance/perfRunner.py @@ -58,8 +58,14 @@ 'f4E2M1FN': 'f32' } # rocmlir-gen host-harness kernel repeat count (--kernel-repeats, used with -ph). +# Also used as the tuning-driver --num-iterations count in legacy benchmark mode. MLIR_N_REPEATS = 100 +# Warmup run count passed to the tuning driver (--warmup-iterations) only in +# legacy benchmark mode; the default do_bench path derives warmup from a time +# budget instead. +WARMUP_ITERATIONS = 10 + # Time budgets (ms) for the tuning-driver benchmark. The number of warmup and # measured iterations is derived from these budgets and the estimated per-launch # runtime (Triton do_bench style). These mirror Triton's do_bench defaults. @@ -69,6 +75,12 @@ TUNE_REP_MS = 100 SLEEP_US = 100 # 0.1 ms +# Sleep between benchmark launches (--sleep-us) used only in legacy benchmark +# mode. This restores the original pre-do_bench perfRunner value (1 ms) so that +# legacy runs reproduce historical timings exactly, independent of the do_bench +# default above. +LEGACY_SLEEP_US = 1000 # 1 ms + FILTER_LAYOUT_MAP = {'N': 'k', 'C': 'c', 'H': 'y', 'W': 'x', 'G': 'g', '0': '0', '1': '1'} INPUT_LAYOUT_MAP = {'N': 'n', 'C': 'c', 'H': 'h', 'W': 'w', 'G': 'g', '0': '0', '1': '1'} OUTPUT_LAYOUT_MAP = {'N': 'n', 'C': 'k', 'H': 'h', 'W': 'w', 'G': 'g', '0': '0', '1': '1'} @@ -2036,15 +2048,18 @@ def run_config_with_mlir(config: PerfConfiguration, if use_tuning_driver: if debug: print("Using HIP timing for benchmarking") + sleep_us = LEGACY_SLEEP_US if legacy_benchmark_mode else SLEEP_US tuning_driver_command = [ paths.mlir_paths.rocmlir_tuning_driver_path, f'--benchmark-config={config.perfconfig}', - f'--rep={TUNE_REP_MS}', f'--warmup={TUNE_WARMUP_MS}', f'--sleep-us={SLEEP_US}', + f'--rep={TUNE_REP_MS}', f'--warmup={TUNE_WARMUP_MS}', f'--sleep-us={sleep_us}', '--use-median' ] if flush_last_level_cache: tuning_driver_command.append("--flush-last-level-cache") if legacy_benchmark_mode: tuning_driver_command.append("--legacy-benchmark-mode") + tuning_driver_command.append(f'--num-iterations={MLIR_N_REPEATS}') + tuning_driver_command.append(f'--warmup-iterations={WARMUP_ITERATIONS}') tuning_driver_command.append('-') outs, noerr = run_pipeline([rocmlir_gen_cmd.split(), tuning_driver_command]) if noerr: diff --git a/mlir/utils/performance/tuningRunner.py b/mlir/utils/performance/tuningRunner.py index 3bbb2494557a..b3bce23b561c 100755 --- a/mlir/utils/performance/tuningRunner.py +++ b/mlir/utils/performance/tuningRunner.py @@ -70,8 +70,20 @@ # ============================================================================= # rocmlir-gen host-harness kernel repeat count (--kernel-repeats, used with -ph). +# Also used as the tuning-driver --num-iterations count in legacy benchmark mode. MLIR_N_REPEATS = 10 +# Warmup run count passed to the tuning driver (--warmup-iterations) only in +# legacy benchmark mode; the default do_bench path derives warmup from a time +# budget instead. +WARMUP_ITERATIONS = 1 + +# Sleep between benchmark launches (--sleep-us) used only in legacy benchmark +# mode. This restores the original pre-do_bench tuningRunner value (0.1 ms) so +# that legacy runs reproduce historical timings exactly, independent of the +# imported do_bench default (SLEEP_US). +LEGACY_SLEEP_US = 100 # 0.1 ms + # A GPU run timeout is different from the outer tuning subprocess timeout: an # in-process kernel may have hung and left the HIP context untrustworthy, so the # tuning driver exits the whole process with this distinct code. @@ -1452,12 +1464,13 @@ def tune_config(test_vector: str, conf_class: type, paths: Paths, options: Optio """Tune a single configuration and return the result.""" gpu_logger = get_gpu_logger(gpu_id) + sleep_us = LEGACY_SLEEP_US if options.legacy_benchmark_mode else SLEEP_US tuning_driver_args = [ f"--tuning-space={options.tuning_space_kind}", f"--rep={TUNE_REP_MS}", f"--warmup={TUNE_WARMUP_MS}", "--use-median", - f"--sleep-us={SLEEP_US}", + f"--sleep-us={sleep_us}", f"--show-all-measurements={options.debug}", f"--num-compile-threads={num_compile_threads}", f"--wait-for-compiles={options.wait_for_compiles}", @@ -1467,6 +1480,8 @@ def tune_config(test_vector: str, conf_class: type, paths: Paths, options: Optio tuning_driver_args.append("--flush-last-level-cache") if options.legacy_benchmark_mode: tuning_driver_args.append("--legacy-benchmark-mode") + tuning_driver_args.append(f"--num-iterations={MLIR_N_REPEATS}") + tuning_driver_args.append(f"--warmup-iterations={WARMUP_ITERATIONS}") env = make_isolated_gpu_env(gpu_id) From bd10009124a0339e3038a4bcd62284da394f03cc Mon Sep 17 00:00:00 2001 From: Daniel Hernandez Date: Tue, 18 Aug 2026 16:45:07 +0200 Subject: [PATCH 6/6] Addressing PR comments --- .../rocmlir-tuning-driver.cpp | 55 ++++++++------- mlir/utils/performance/perfRunner.py | 67 ++++++++++--------- mlir/utils/performance/tuningRunner.py | 32 ++++----- 3 files changed, 84 insertions(+), 70 deletions(-) diff --git a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp index c0e028b4f09a..6eb13fd04121 100644 --- a/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp +++ b/mlir/tools/rocmlir-tuning-driver/rocmlir-tuning-driver.cpp @@ -137,37 +137,40 @@ static llvm::cl::opt rep( "rep", llvm::cl::desc("Target benchmark time in milliseconds. The number of " "measured iterations is derived from this budget and the " - "estimated per-launch runtime (Triton do_bench style)."), + "estimated per-launch runtime (Triton do_bench style). Only " + "used with --triton-benchmark-mode."), llvm::cl::value_desc("benchmark milliseconds"), llvm::cl::init(100)); static llvm::cl::opt warmup( "warmup", llvm::cl::desc("Target warmup time in milliseconds. The number of warmup " "iterations is derived from this budget and the estimated " - "per-launch runtime (Triton do_bench style)."), + "per-launch runtime (Triton do_bench style). Only used with " + "--triton-benchmark-mode."), llvm::cl::value_desc("warmup milliseconds"), llvm::cl::init(25)); -static llvm::cl::opt legacyBenchmarkMode( - "legacy-benchmark-mode", +static llvm::cl::opt tritonBenchmarkMode( + "triton-benchmark-mode", llvm::cl::desc( - "Use the legacy rocMLIR benchmarking method (fixed iteration counts " - "with a small-vs-large-kernel split) instead of the default " - "Triton do_bench-style time-budget measurement. Kept for " - "apples-to-apples comparison against older rocMLIR versions. In this " - "mode --num-iterations/--warmup-iterations control the run counts and " - "--rep/--warmup are ignored."), + "Use the Triton do_bench-style time-budget measurement (iteration " + "counts derived from --rep/--warmup and the estimated per-launch " + "runtime) instead of the default rocMLIR benchmarking method. " + "Enable this for apples-to-apples comparison against Triton. In the " + "default mode --num-iterations/--warmup-iterations control the " + "run counts and --rep/--warmup are ignored."), llvm::cl::init(false)); static llvm::cl::opt numIterations( "num-iterations", llvm::cl::desc("Number of times to run each kernel for averaging (only " - "used with --legacy-benchmark-mode)"), + "used in the default mode, i.e. when " + "--triton-benchmark-mode is not set)"), llvm::cl::value_desc("number of runs"), llvm::cl::init(100)); static llvm::cl::opt warmupIterations( "warmup-iterations", - llvm::cl::desc( - "Number of warmup runs (only used with --legacy-benchmark-mode)"), + llvm::cl::desc("Number of warmup runs (only used in the default mode, i.e. " + "when --triton-benchmark-mode is not set)"), llvm::cl::value_desc("number of warmup runs"), llvm::cl::init(10)); static llvm::cl::opt @@ -373,10 +376,12 @@ struct BenchmarkParams { bool waitForCompiles; unsigned gpuRunTimeoutSec; bool flushLastLevelCache; - // Legacy benchmarking path (fixed iteration counts + small/large kernel - // split). When false, the default Triton do_bench-style measurement is used - // and numIterations/warmupIterations are ignored. - bool legacyBenchmarkMode; + // When true, use the Triton do_bench-style time-budget measurement (deriving + // iteration counts from rep/warmup) and ignore + // numIterations/warmupIterations. When false (the default) the legacy + // benchmarking path is used (fixed iteration counts + small/large kernel + // split). + bool tritonBenchmarkMode; unsigned numIterations; unsigned warmupIterations; }; @@ -449,7 +454,8 @@ struct ThreadResources { }; // Legacy small-kernel path: run all iterations back-to-back and time the whole -// batch with a single CPU timer. Only used with --legacy-benchmark-mode. +// batch with a single CPU timer. Used in the default mode (i.e. when +// --triton-benchmark-mode is not set). static LogicalResult measureSmallKernel(unsigned iterations, hipStream_t stream, const std::vector &functions, @@ -492,7 +498,8 @@ measureSmallKernel(unsigned iterations, hipStream_t stream, } // Legacy large-kernel path: time each iteration individually with GPU events, -// synchronizing after every launch. Only used with --legacy-benchmark-mode. +// synchronizing after every launch. Used in the default mode (i.e. when +// --triton-benchmark-mode is not set). static LogicalResult measureLargeKernel( unsigned iterations, hipStream_t stream, const std::vector &functions, ArrayRef blockSizes, @@ -660,10 +667,10 @@ benchmarkKernels(ArrayRef binaries, bool isSmallKernel = false; unsigned smallKernelIters = 0; - if (params.legacyBenchmarkMode) { - // Legacy benchmarking: fixed --num-iterations/--warmup-iterations counts - // with a small-vs-large-kernel split. Kept for apples-to-apples comparison - // with older rocMLIR versions. + if (!params.tritonBenchmarkMode) { + // Legacy benchmarking (default): fixed --num-iterations/--warmup-iterations + // counts with a small-vs-large-kernel split. Kept as the default for + // apples-to-apples comparison with older rocMLIR versions. bool benchmarkMode = !params.benchmarkConfig.empty(); unsigned iterations = params.numIterations; @@ -1013,7 +1020,7 @@ static LogicalResult runTuningLoop(ModuleOp source) { waitForCompiles, gpuRunTimeout, flushLastLevelCache, - legacyBenchmarkMode, + tritonBenchmarkMode, numIterations, warmupIterations}; diff --git a/mlir/utils/performance/perfRunner.py b/mlir/utils/performance/perfRunner.py index 2c88748fb2ab..60b94cb94224 100644 --- a/mlir/utils/performance/perfRunner.py +++ b/mlir/utils/performance/perfRunner.py @@ -58,27 +58,31 @@ 'f4E2M1FN': 'f32' } # rocmlir-gen host-harness kernel repeat count (--kernel-repeats, used with -ph). -# Also used as the tuning-driver --num-iterations count in legacy benchmark mode. +# Also used as the tuning-driver --num-iterations count in the default benchmark +# mode. MLIR_N_REPEATS = 100 -# Warmup run count passed to the tuning driver (--warmup-iterations) only in -# legacy benchmark mode; the default do_bench path derives warmup from a time -# budget instead. +# Warmup run count passed to the tuning driver (--warmup-iterations) in the +# default benchmark mode; the opt-in Triton do_bench path derives warmup from a +# time budget instead. WARMUP_ITERATIONS = 10 -# Time budgets (ms) for the tuning-driver benchmark. The number of warmup and +# Time budgets (ms) for the tuning-driver benchmark, used only with the opt-in +# Triton do_bench path (--triton-benchmark-mode). The number of warmup and # measured iterations is derived from these budgets and the estimated per-launch -# runtime (Triton do_bench style). These mirror Triton's do_bench defaults. -# tuningRunner imports these so that inference (benchmark) numbers stay consistent -# with tuning numbers for the same perfConfig. +# runtime. These mirror Triton's do_bench defaults. tuningRunner imports these so +# that inference (benchmark) numbers stay consistent with tuning numbers for the +# same perfConfig. TUNE_WARMUP_MS = 25 TUNE_REP_MS = 100 +# Sleep between benchmark launches (--sleep-us) used with the opt-in Triton +# do_bench path (--triton-benchmark-mode). SLEEP_US = 100 # 0.1 ms -# Sleep between benchmark launches (--sleep-us) used only in legacy benchmark -# mode. This restores the original pre-do_bench perfRunner value (1 ms) so that -# legacy runs reproduce historical timings exactly, independent of the do_bench -# default above. +# Sleep between benchmark launches (--sleep-us) used in the default benchmark +# mode. This preserves the original pre-do_bench perfRunner value (1 ms) so that +# default-mode runs reproduce historical timings exactly, independent of the +# Triton do_bench value above. LEGACY_SLEEP_US = 1000 # 1 ms FILTER_LAYOUT_MAP = {'N': 'k', 'C': 'c', 'H': 'y', 'W': 'x', 'G': 'g', '0': '0', '1': '1'} @@ -2024,7 +2028,7 @@ def run_config_with_mlir(config: PerfConfiguration, rocmlir_gen_flags, use_rocprof=False, flush_last_level_cache=False, - legacy_benchmark_mode=False, + triton_benchmark_mode=False, debug=True): # remove the result file generated by rocprof in previous benchmarking if os.path.exists(get_profiler_output_path(arch, BENCHMARKING_STATS_FILE_NAME)): @@ -2048,7 +2052,7 @@ def run_config_with_mlir(config: PerfConfiguration, if use_tuning_driver: if debug: print("Using HIP timing for benchmarking") - sleep_us = LEGACY_SLEEP_US if legacy_benchmark_mode else SLEEP_US + sleep_us = SLEEP_US if triton_benchmark_mode else LEGACY_SLEEP_US tuning_driver_command = [ paths.mlir_paths.rocmlir_tuning_driver_path, f'--benchmark-config={config.perfconfig}', f'--rep={TUNE_REP_MS}', f'--warmup={TUNE_WARMUP_MS}', f'--sleep-us={sleep_us}', @@ -2056,8 +2060,9 @@ def run_config_with_mlir(config: PerfConfiguration, ] if flush_last_level_cache: tuning_driver_command.append("--flush-last-level-cache") - if legacy_benchmark_mode: - tuning_driver_command.append("--legacy-benchmark-mode") + if triton_benchmark_mode: + tuning_driver_command.append("--triton-benchmark-mode") + else: tuning_driver_command.append(f'--num-iterations={MLIR_N_REPEATS}') tuning_driver_command.append(f'--warmup-iterations={WARMUP_ITERATIONS}') tuning_driver_command.append('-') @@ -2102,7 +2107,7 @@ def benchmark_mlir(commandline, rocmlir_gen_flags, use_rocprof=False, flush_last_level_cache=False, - legacy_benchmark_mode=False): + triton_benchmark_mode=False): config = conf_class.from_command_line(commandline, arch, num_cu, num_chiplets) config_str = config.to_command_line() if tuning_db: @@ -2112,7 +2117,7 @@ def benchmark_mlir(commandline, return config.table_entry(np.nan) nanoseconds = run_config_with_mlir(config, paths, arch, rocmlir_gen_flags, use_rocprof, - flush_last_level_cache, legacy_benchmark_mode) + flush_last_level_cache, triton_benchmark_mode) return config.table_entry(nanoseconds) @@ -2128,26 +2133,26 @@ def generate_performance_results(configs, rocmlir_gen_flags, use_rocprof=False, flush_last_level_cache=False, - legacy_benchmark_mode=False): + triton_benchmark_mode=False): # Never pass tuning DB to this run mlir_df = pd.DataFrame( benchmark_mlir(test_vector.split( sep=' '), conf_class, paths, arch, num_cu, num_chiplets, None, rocmlir_gen_flags, - use_rocprof, flush_last_level_cache, legacy_benchmark_mode) + use_rocprof, flush_last_level_cache, triton_benchmark_mode) for test_vector in configs) tuned_df = None if tuning_db: tuned_df = pd.DataFrame( benchmark_mlir(test_vector.split(sep=' '), conf_class, paths, arch, num_cu, num_chiplets, tuning_db, rocmlir_gen_flags, use_rocprof, - flush_last_level_cache, legacy_benchmark_mode) + flush_last_level_cache, triton_benchmark_mode) for test_vector in configs) quick_tuned_df = None if quick_tuning_db: quick_tuned_df = pd.DataFrame( benchmark_mlir(test_vector.split(sep=' '), conf_class, paths, arch, num_cu, num_chiplets, quick_tuning_db, rocmlir_gen_flags, use_rocprof, - flush_last_level_cache, legacy_benchmark_mode) + flush_last_level_cache, triton_benchmark_mode) for test_vector in configs) external_df = pd.DataFrame( @@ -2392,7 +2397,7 @@ def benchmark_fusion_kernels(test_dir, tuning_db: MaybeTuningDb, use_rocprof=False, flush_last_level_cache=False, - legacy_benchmark_mode=False): + triton_benchmark_mode=False): all_tests = [] # filename, test_vector, fut_name perf_results = {} # associate test_vector to config and performances chip = GFX_CHIP_RE.search(arch).group(0) @@ -2468,7 +2473,7 @@ def benchmark_fusion_kernels(test_dir, # Run gemm or conv op with the same configuration nanoseconds = run_config_with_mlir(config, paths, arch, '', use_rocprof, - flush_last_level_cache, legacy_benchmark_mode) + flush_last_level_cache, triton_benchmark_mode) one_entry['MLIR TFlops'] = config.compute_tflops(nanoseconds) one_entry['Fusion/MLIR'] = one_entry['TFlops'] / one_entry['MLIR TFlops'] one_entry['FileName'] = filename @@ -2707,11 +2712,11 @@ def main(args=None): ) parser.add_argument( - "--legacy-benchmark-mode", + "--triton-benchmark-mode", action='store_true', default=False, help= - "Use the legacy rocMLIR benchmarking method (fixed iteration counts with a small-vs-large-kernel split) instead of the default Triton do_bench-style time-budget measurement. Kept for apples-to-apples comparison against older rocMLIR versions." + "Use the Triton do_bench-style time-budget measurement (iteration counts derived from time budgets) instead of the default rocMLIR benchmarking method (fixed iteration counts with a small-vs-large-kernel split). Enable this for apples-to-apples comparison against Triton." ) parsed_args = parser.parse_args(args) @@ -2788,7 +2793,7 @@ def main(args=None): generate_performance_results(configs, conf_class, paths, arch, num_cu, num_chiplets, tuning_db, quick_tuning_db, rocmlir_gen_flags, parsed_args.use_rocprof, parsed_args.flush_last_level_cache, - parsed_args.legacy_benchmark_mode) + parsed_args.triton_benchmark_mode) elif parsed_args.tuning: tune_mlir_kernels(configs, arch, num_cu, num_chiplets) elif optype == Operation.FUSION: @@ -2798,14 +2803,14 @@ def main(args=None): benchmark_fusion_kernels(parsed_args.test_dir, paths, arch, num_cu, num_chiplets, tuning_db, parsed_args.use_rocprof, parsed_args.flush_last_level_cache, - parsed_args.legacy_benchmark_mode) + parsed_args.triton_benchmark_mode) else: if parsed_args.batch_mlir: df = pd.DataFrame( benchmark_mlir(test_vector.split( sep=' '), conf_class, paths, arch, num_cu, num_chiplets, tuning_db, rocmlir_gen_flags, parsed_args.use_rocprof, parsed_args. - flush_last_level_cache, parsed_args.legacy_benchmark_mode) + flush_last_level_cache, parsed_args.triton_benchmark_mode) for test_vector in configs) elif parsed_args.batch_external: df = pd.DataFrame( @@ -2826,14 +2831,14 @@ def main(args=None): benchmark_mlir(parsed_args.config, conf_class, paths, arch, num_cu, num_chiplets, tuning_db, rocmlir_gen_flags, parsed_args.use_rocprof, parsed_args.flush_last_level_cache, - parsed_args.legacy_benchmark_mode) + parsed_args.triton_benchmark_mode) ]) else: df = pd.DataFrame([ benchmark_mlir(config.split(), conf_class, paths, arch, num_cu, num_chiplets, tuning_db, rocmlir_gen_flags, parsed_args.use_rocprof, parsed_args.flush_last_level_cache, - parsed_args.legacy_benchmark_mode) for config in configs + parsed_args.triton_benchmark_mode) for config in configs ]) df.to_csv(parsed_args.filename) with pd.option_context('display.precision', reportUtils.ROUND_DIGITS): diff --git a/mlir/utils/performance/tuningRunner.py b/mlir/utils/performance/tuningRunner.py index b3bce23b561c..51334aa2490d 100755 --- a/mlir/utils/performance/tuningRunner.py +++ b/mlir/utils/performance/tuningRunner.py @@ -70,18 +70,19 @@ # ============================================================================= # rocmlir-gen host-harness kernel repeat count (--kernel-repeats, used with -ph). -# Also used as the tuning-driver --num-iterations count in legacy benchmark mode. +# Also used as the tuning-driver --num-iterations count in the default benchmark +# mode. MLIR_N_REPEATS = 10 -# Warmup run count passed to the tuning driver (--warmup-iterations) only in -# legacy benchmark mode; the default do_bench path derives warmup from a time -# budget instead. +# Warmup run count passed to the tuning driver (--warmup-iterations) in the +# default benchmark mode; the opt-in Triton do_bench path derives warmup from a +# time budget instead. WARMUP_ITERATIONS = 1 -# Sleep between benchmark launches (--sleep-us) used only in legacy benchmark -# mode. This restores the original pre-do_bench tuningRunner value (0.1 ms) so -# that legacy runs reproduce historical timings exactly, independent of the -# imported do_bench default (SLEEP_US). +# Sleep between benchmark launches (--sleep-us) used in the default benchmark +# mode. This preserves the original pre-do_bench tuningRunner value (0.1 ms) so +# that default-mode runs reproduce historical timings exactly, independent of the +# imported Triton do_bench value (SLEEP_US). LEGACY_SLEEP_US = 100 # 0.1 ms # A GPU run timeout is different from the outer tuning subprocess timeout: an @@ -211,7 +212,7 @@ class Options: num_cpus: Optional[int] wait_for_compiles: bool flush_last_level_cache: bool - legacy_benchmark_mode: bool + triton_benchmark_mode: bool timeout: Optional[int] verify_timeout: Optional[int] gpu_run_timeout: int @@ -1464,7 +1465,7 @@ def tune_config(test_vector: str, conf_class: type, paths: Paths, options: Optio """Tune a single configuration and return the result.""" gpu_logger = get_gpu_logger(gpu_id) - sleep_us = LEGACY_SLEEP_US if options.legacy_benchmark_mode else SLEEP_US + sleep_us = SLEEP_US if options.triton_benchmark_mode else LEGACY_SLEEP_US tuning_driver_args = [ f"--tuning-space={options.tuning_space_kind}", f"--rep={TUNE_REP_MS}", @@ -1478,8 +1479,9 @@ def tune_config(test_vector: str, conf_class: type, paths: Paths, options: Optio ] if options.flush_last_level_cache: tuning_driver_args.append("--flush-last-level-cache") - if options.legacy_benchmark_mode: - tuning_driver_args.append("--legacy-benchmark-mode") + if options.triton_benchmark_mode: + tuning_driver_args.append("--triton-benchmark-mode") + else: tuning_driver_args.append(f"--num-iterations={MLIR_N_REPEATS}") tuning_driver_args.append(f"--warmup-iterations={WARMUP_ITERATIONS}") @@ -2166,11 +2168,11 @@ def parse_arguments(gpu_topology: GpuTopology, ) parser.add_argument( - "--legacy-benchmark-mode", + "--triton-benchmark-mode", action='store_true', default=False, help= - "Use the legacy rocMLIR benchmarking method (fixed iteration counts with a small-vs-large-kernel split) instead of the default Triton do_bench-style time-budget measurement. Kept for apples-to-apples comparison against older rocMLIR versions." + "Use the Triton do_bench-style time-budget measurement (iteration counts derived from time budgets) instead of the default rocMLIR benchmarking method (fixed iteration counts with a small-vs-large-kernel split). Enable this for apples-to-apples comparison against Triton." ) parser.add_argument("-s", @@ -2276,7 +2278,7 @@ def main(args=None): num_cpus=parsed_args.num_cpus, wait_for_compiles=parsed_args.wait_for_compiles, flush_last_level_cache=parsed_args.flush_last_level_cache, - legacy_benchmark_mode=parsed_args.legacy_benchmark_mode, + triton_benchmark_mode=parsed_args.triton_benchmark_mode, timeout=parsed_args.timeout, verify_timeout=parsed_args.verify_timeout, gpu_run_timeout=parsed_args.gpu_run_timeout)