From 669bd62b1f80a821643f5a9ac43040bc80d2386b Mon Sep 17 00:00:00 2001 From: Brad Pepers Date: Sat, 22 Aug 2026 22:44:04 -0600 Subject: [PATCH 1/4] fix(miopen): skip benchmarking oversized naive conv during Find to avoid TDR MIOpen's Find step times every applicable solver by launching it. The naive convolution kernel is un-tiled, so on a large problem a single launch runs for seconds and merely benchmarking it trips the OS GPU watchdog (a TDR / driver reset) even when a fast solver also applies and would win. Skip *running* naive in Find's mini-benchmark when both a non-naive solver is applicable (computed from the full applicable list before workspace filtering, closing the leak where the alternative was workspace-filtered and naive ran anyway) and the problem's total MAC work exceeds a threshold (ConvDirectNaiveConvExceedsWorkLimit, ~16 GMAC, overridable via MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK). Naive stays applicable at every size, so it still runs when it is the only option, and small problems where it is fastest (including all real depthwise convs) keep competing. Consolidates the naive-skip policy into one ShouldSkipNaiveBenchmark predicate, drops the reorder/out-parameter threading in EvaluateInvokers, and removes the superseded output-size-based MIOPEN_CONV_DIRECT_MAX_SIZE gate. Adds a CPU-only gtest for the work metric (no Handle, no kernel launch). Relates to #3509, #1860. Co-Authored-By: Claude Opus 4.8 --- projects/miopen/CHANGELOG.md | 4 + projects/miopen/src/conv/solver_finders.cpp | 151 +++++++++---- projects/miopen/src/env_debug.cpp | 3 + .../include/miopen/conv/solver_finders.hpp | 4 +- .../miopen/solver/conv_direct_naive_conv.hpp | 11 + .../solver/conv/conv_direct_naive_conv.cpp | 53 +++++ .../unit_conv_solver_naive_work_gate.cpp | 199 ++++++++++++++++++ 7 files changed, 387 insertions(+), 38 deletions(-) create mode 100644 projects/miopen/test/gtest/unit_conv_solver_naive_work_gate.cpp diff --git a/projects/miopen/CHANGELOG.md b/projects/miopen/CHANGELOG.md index b309592784e9..3e6b9d6dc594 100644 --- a/projects/miopen/CHANGELOG.md +++ b/projects/miopen/CHANGELOG.md @@ -9,7 +9,11 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj * [Conv] Added gfx950 depthwise backward-weights (fp16/bf16) and gfx1250 depthwise kernels to the `ConvHipConv` solver (hipconv v0.3.1). +### Resolved Issues +* [Conv] Fixed a GPU driver watchdog reset (TDR) that could occur while find benchmarked the un-tiled naive convolution solver on very large problems (e.g. VAE-decode and 3D convolutions). Naive is now skipped during find when a non-naive solver is applicable and the problem's total MAC work exceeds a threshold (~16 GMAC, overridable via `MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK`). This complements the existing `MIOPEN_NAIVE_TIMEOUT` deferral, which bounds how long find waits on the naive warmup but cannot prevent the launch itself. Naive still runs when it is the only applicable solver, and small problems where it is fastest continue to compete. + ### Removed +* [Conv] Removed the `MIOPEN_CONV_DIRECT_MAX_SIZE` environment variable (an output-element-count threshold that disabled the entire Direct algorithm). It was a poor proxy for kernel launch cost and is superseded by the naive-conv work gate described above. * [Conv] Removed gfx803 convolution solver `ConvBinWinogradRxSFused` and its kernel sources. * [Conv] Removed the gfx803 code paths and kernel sources from `ConvBinWinograd3x3U` and `ConvBinWinogradRxS`. * [Conv] Removed the gfx803 find-database files. diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index 88036aaf17e1..fd8b5538c227 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -21,7 +21,9 @@ #include #include #include +#include #include +#include #include MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_GEMM) @@ -38,7 +40,6 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_NAIVE_TIMEOUT, true) MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_NAIVE_TIMEOUT_FACTOR, 300) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_SEARCH_CUTOFF, false) MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_FIND_SKIP_PCT, 130) -MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_CONV_DIRECT_MAX_SIZE, 0) namespace miopen { @@ -352,6 +353,74 @@ static NaiveWarmup TryNaiveWithTimeout(const Handle& handle, return {NaiveWarmup::Status::TimedOut, 0.0f}; } +/// A solver is Naive iff its id contains "Naive" (ConvDirectNaiveConv{Fwd,Bwd,Wrw}). +/// Single source of truth for the string test used by the per-solver skip check, the +/// timeout-deferral ordering, and the FindCore "does any non-Naive solver apply" scan. +static bool IsNaiveSolverId(const std::string& solver_id) +{ + return solver_id.find("Naive") != std::string::npos; +} + +/// Decide whether the Naive convolution solver should be *skipped* -- i.e. not +/// executed in Find's mini-benchmark -- for the current solution. +/// +/// Why this exists: Naive is MIOpen's universal fallback and stays *applicable* at +/// any problem size, but its kernel is un-tiled, so on a large problem a single +/// launch runs for multiple seconds. Find times every applicable solver by actually +/// *executing* it, so merely benchmarking Naive on such a shape trips the OS GPU +/// watchdog (a TDR / driver reset) even when a fast solver (CK/GEMM) also applies +/// and would ultimately win. Naive does not have to be *selected* to cause the +/// hang -- being *benchmarked* is enough. +/// +/// This is deliberately a *pre-launch* gate, complementing (not replacing) the +/// TryNaiveWithTimeout path below. That path bounds how long Find *waits* on the +/// naive warmup, but the dispatch has already happened and the kernel keeps +/// occupying the GPU after the slot is abandoned -- which is exactly what the OS +/// watchdog measures, so abandoning the wait does not avert a TDR. It also only +/// engages once a non-Naive solver has actually *succeeded* (non_naive_succeeded), +/// so it does nothing when every alternative is rejected at evaluation -- the leak +/// this gate closes by keying off applicability instead. We therefore skip +/// *launching* Naive during the benchmark exactly when BOTH of these hold: +/// +/// * non_naive_exists -- some non-Naive solver was found applicable for this +/// problem. The caller computes this from the full +/// solution list *before* workspace filtering, so a fast +/// solver that is later workspace-filtered still counts; +/// that is what closes the TDR leak where the alternative +/// never gets to time itself. AND +/// * naive_exceeds_work -- the problem's total MAC work is over the Naive work +/// limit (~16 GMAC, see ConvDirectNaiveConvExceedsWorkLimit), +/// i.e. large enough that a single launch could TDR. +/// +/// Consequences of this exact condition (the whole selection policy in one place): +/// * Naive is the *sole* applicable solver -> not skipped -> it runs, preserving +/// the universal-fallback guarantee. A huge sole-Naive shape may then still TDR +/// -- an honest "extend coverage here" signal we deliberately do not mask. +/// * Any shape under the work limit -> not skipped -> Naive competes and wins on +/// merit wherever it is fastest. This notably covers *all* real depthwise convs +/// (group == C == K => C_per_group == 1 => ~C x less work), which never approach +/// the limit and for which Naive is often the fastest option (e.g. NCHW): they +/// keep competing with no special case. +/// * Only large-and-avoidable shapes (over the limit with an alternative present) +/// are skipped -- the actual TDR case. +/// +/// MIOPEN_NAIVE_TIMEOUT (naive_timeout, default on) is the master switch for Naive +/// deferral and gates this policy as well as the timeout path; unset it to force +/// Naive to always compete. The work limit itself is separately tunable via +/// MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK (set it very large to disable just this +/// gate while keeping the timeout). Golden references are unaffected: GpuConvReference +/// compiles and launches the naive kernel directly, bypassing the solver framework, so +/// verification never reaches this gate. +static bool ShouldSkipNaiveBenchmark(bool is_naive, + bool naive_timeout, + bool non_naive_exists, + bool naive_exceeds_work) +{ + if(!is_naive || !naive_timeout) + return false; + return non_naive_exists && naive_exceeds_work; +} + /// Register invoker only for the best solution within algorithm. std::vector EvaluateInvokers(const Handle& handle, const std::vector& solutions, @@ -360,7 +429,9 @@ std::vector EvaluateInvokers(const Handle& handle, const AnyInvokeParams& invoke_ctx, FindCoreResult& core_result, bool force_attach_binary, - bool& non_naive_succeeded) + bool& non_naive_succeeded, + bool non_naive_exists, + bool naive_exceeds_work) { std::vector ret; @@ -369,7 +440,7 @@ std::vector EvaluateInvokers(const Handle& handle, return ret; const auto is_naive_solver = [](const solver::ConvSolution& s) { - return s.solver_id.find("Naive") != std::string::npos; + return IsNaiveSolverId(s.solver_id); }; bool naive_timeout = env::value(MIOPEN_NAIVE_TIMEOUT); @@ -399,7 +470,24 @@ std::vector EvaluateInvokers(const Handle& handle, { const auto& sol = solutions[idx]; - const bool is_naive = is_naive_solver(sol); + const bool is_naive = is_naive_solver(sol); + if(ShouldSkipNaiveBenchmark(is_naive, naive_timeout, non_naive_exists, naive_exceeds_work)) + { + MIOPEN_LOG_I("Skipping Naive Solver: " << algorithm_name.ToString() << ":" + << sol.solver_id); + continue; + } + if(naive_timeout && is_naive) + { + // Naive is being kept in the benchmark set even though the work gate is on. + // Name the reason so this reads as an expected retention, not an anomaly. It may + // still be cut short by the timeout path below (cutoff_naive). + const auto* reason = !non_naive_exists ? "sole applicable solver" + : !naive_exceeds_work ? "below work limit" + : "skip criteria not met"; + MIOPEN_LOG_I("Retaining Naive Solver (" << reason << "): " << algorithm_name.ToString() + << ":" << sol.solver_id); + } const bool cutoff_naive = defer_naive && is_naive && non_naive_succeeded; if(!conv::IsEnoughWorkspace( @@ -635,6 +723,23 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, ret.solutions.reserve(total); + // Does any non-Naive solver apply, across all algorithms? Computed from the full solution + // list *before* per-solver workspace filtering, so a fast-but-workspace-limited alternative + // still counts and Naive is deferred rather than benchmarked (the TDR-leak fix). Also decide + // once whether this conv's total MAC work is large enough that the un-tiled Naive kernel would + // trip the OS GPU watchdog if benchmarked. Non-conv (fusion) problems yield nullptr and are + // inert here. + const bool non_naive_exists = + std::any_of(solutions.begin(), solutions.end(), [](const auto& g) { + return std::any_of(g.second.begin(), g.second.end(), [](const auto& s) { + return !IsNaiveSolverId(s.solver_id); + }); + }); + const auto* conv_problem = dynamic_cast(&problem); + const bool naive_exceeds_work = + (conv_problem != nullptr) && + solver::conv::ConvDirectNaiveConvExceedsWorkLimit(*conv_problem); + bool non_naive_succeeded = false; for(const auto& ss : solutions) { @@ -645,7 +750,9 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, invoke_ctx, ret, force_attach_binary, - non_naive_succeeded); + non_naive_succeeded, + non_naive_exists, + naive_exceeds_work); ret.solutions.insert(ret.solutions.end(), std::make_move_iterator(evaluated.begin()), @@ -657,37 +764,7 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, namespace conv { -namespace detail { -/// Determine if problem size exceeds threshold for Direct solver. -/// -/// The result tensor is used to estimate problem size. -/// The maximum size is determined by MIOPEN_CONV_DIRECT_MAX_SIZE environment variable. -/// -/// @param problem The convolution problem description. -bool IsDirectProblemTooLarge(const ProblemDescription& problem) -{ - const unsigned long long max_size = env::value(MIOPEN_CONV_DIRECT_MAX_SIZE); - // 0 means no limit - if(max_size == 0) - return false; - - // For FWD/BWD: 'out' is the result (swapped in BWD) - // For WRW: 'weights' is the result (out is dy, not dw) - const size_t problem_size = problem.IsDirectionBackwardWrW() - ? problem.GetWeights().GetElementSize() - : problem.GetOut().GetElementSize(); - - // Problem size is within limit - if(problem_size <= max_size) - return false; - - MIOPEN_LOG_I2("DirectSolverFinder disabled for problem size " - << problem_size << " > " << max_size << " (MIOPEN_CONV_DIRECT_MAX_SIZE)"); - return true; -} -} // namespace detail - -bool IsAlgorithmDisabled(miopenConvAlgorithm_t algo, const ProblemDescription& problem) +bool IsAlgorithmDisabled(miopenConvAlgorithm_t algo, const ProblemDescription& /*problem*/) { switch(algo) { // clang-format off @@ -698,7 +775,7 @@ bool IsAlgorithmDisabled(miopenConvAlgorithm_t algo, const ProblemDescription& p return true; #endif case miopenConvolutionAlgoDirect: - return env::disabled(MIOPEN_DEBUG_CONV_DIRECT) || detail::IsDirectProblemTooLarge(problem); + return env::disabled(MIOPEN_DEBUG_CONV_DIRECT); case miopenConvolutionAlgoFFT: return env::disabled(MIOPEN_DEBUG_CONV_FFT); case miopenConvolutionAlgoWinograd: diff --git a/projects/miopen/src/env_debug.cpp b/projects/miopen/src/env_debug.cpp index 317d86221c9b..3c5cca59e3a9 100644 --- a/projects/miopen/src/env_debug.cpp +++ b/projects/miopen/src/env_debug.cpp @@ -50,6 +50,7 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3) MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_DIRECT) +MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_FFT) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_GEMM) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM) @@ -193,6 +194,8 @@ const LibEnvVar& FindEnvVariable(std::string_view name) {MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3.GetName(), MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3}, {MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS.GetName(), MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS}, {MIOPEN_DEBUG_CONV_DIRECT.GetName(), MIOPEN_DEBUG_CONV_DIRECT}, + {MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK.GetName(), + MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK}, {MIOPEN_DEBUG_CONV_FFT.GetName(), MIOPEN_DEBUG_CONV_FFT}, {MIOPEN_DEBUG_CONV_GEMM.GetName(), MIOPEN_DEBUG_CONV_GEMM}, {MIOPEN_DEBUG_CONV_IMPLICIT_GEMM.GetName(), MIOPEN_DEBUG_CONV_IMPLICIT_GEMM}, diff --git a/projects/miopen/src/include/miopen/conv/solver_finders.hpp b/projects/miopen/src/include/miopen/conv/solver_finders.hpp index db4e2141649d..ff0466977e9f 100644 --- a/projects/miopen/src/include/miopen/conv/solver_finders.hpp +++ b/projects/miopen/src/include/miopen/conv/solver_finders.hpp @@ -171,7 +171,9 @@ std::vector EvaluateInvokers(const Handle& handle, const AnyInvokeParams& invoke_ctx, FindCoreResult& core_result, bool force_attach_binary, - bool& non_naive_succeeded); + bool& non_naive_succeeded, + bool non_naive_exists = false, + bool naive_exceeds_work = false); FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, const ExecutionContext& ctx, diff --git a/projects/miopen/src/include/miopen/solver/conv_direct_naive_conv.hpp b/projects/miopen/src/include/miopen/solver/conv_direct_naive_conv.hpp index 202d92786afe..3a15523dd450 100644 --- a/projects/miopen/src/include/miopen/solver/conv_direct_naive_conv.hpp +++ b/projects/miopen/src/include/miopen/solver/conv_direct_naive_conv.hpp @@ -52,6 +52,17 @@ std::string ConvDirectNaiveConvCompileOption(const ExecutionContext& ctx, bool ConvDirectNaiveConvIsApplicableByKernelType(const ExecutionContext&, const miopen::conv::ProblemDescription&); +/// Returns true when the problem's total MAC work exceeds the naive-conv work +/// limit. Naive stays applicable at any size (it is the universal fallback); this +/// is consumed at the selection layer (EvaluateInvokers) to keep the un-tiled +/// naive kernel from being benchmarked when it would run long enough to trip the +/// OS GPU watchdog / TDR and a non-naive alternative also applies. +/// Overridable via MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK; see the .cpp for detail. +/// Exported (MIOPEN_INTERNALS_EXPORT) so the gtest unit test can link against it +/// directly and exercise the metric on the CPU without launching a kernel. +MIOPEN_INTERNALS_EXPORT bool +ConvDirectNaiveConvExceedsWorkLimit(const miopen::conv::ProblemDescription&); + bool IsInputFp32(const miopen::conv::ProblemDescription&); bool IsInputFp16(const miopen::conv::ProblemDescription&); bool IsInputBfp16(const miopen::conv::ProblemDescription&); diff --git a/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp b/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp index 8025014dcfb1..89a2c779a296 100644 --- a/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp +++ b/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp @@ -37,6 +37,11 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_USE_PACKED_KERNELS); +// Runtime override (in MACs) for the naive-conv work-size applicability gate. +// 0 (unset) => use the built-in NAIVE_CONV_MAX_WORK default; set to a very large +// value to effectively disable the gate. See ConvDirectNaiveConvExceedsWorkLimit. +MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK); + namespace miopen { namespace solver { @@ -310,6 +315,54 @@ bool ConvDirectNaiveConvIsApplicableByKernelType(const ExecutionContext& ctx, return true; } +// Upper bound (in multiply-accumulate operations) on the total work above which the +// naive convolution solver should not be *benchmarked* during Find. The naive +// kernel is un-tiled and the GPU saturates, so its wall-time is proportional to the +// total MAC count: N * K * C_per_group * output_spatial_volume * filter_volume. On +// very large problems (e.g. VAE-decode / 3D shapes) a single naive launch runs for +// multiple seconds; when MIOpen's Find step *benchmarks* candidate solvers it will +// execute naive and trip the OS GPU watchdog (TDR) -- even when a fast solver also +// applies and would ultimately win. Naive stays *applicable* at any size so it can +// serve as the universal fallback; the selection layer (EvaluateInvokers) consumes +// this threshold to skip benchmarking naive when it is over the limit AND a +// non-naive alternative also applies. When naive is the sole applicable solver it +// still runs, so a shape only naive can serve is served (and, if huge, may TDR on a +// short-watchdog OS -- an honest "extend coverage here" signal, not masked). +// Overridable via MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK. +// NOTE: the 16 GMAC default is a conservative estimate, not a measured value; it +// should be calibrated once on stable hardware (see project_naive_find_tdr). +constexpr size_t NAIVE_CONV_MAX_WORK = static_cast(16) * 1000 * 1000 * 1000; // ~16 GMAC + +bool ConvDirectNaiveConvExceedsWorkLimit(const ProblemDescription& problem) +{ + const size_t n = problem.GetInBatchSize(); + const size_t k = problem.GetOutChannels(); + const auto group = static_cast(problem.GetGroupCount()); + const size_t c_per_group = (group == 0) ? problem.GetInChannels() // + : problem.GetInChannels() / group; + + const size_t out_vol = problem.GetOutWidth() * problem.GetOutHeight() * + (problem.Is2d() ? size_t{1} : problem.GetOutDepth()); + const size_t fil_vol = problem.GetWeightsWidth() * problem.GetWeightsHeight() * + (problem.Is2d() ? size_t{1} : problem.GetWeightsDepth()); + + const size_t work = n * k * c_per_group * out_vol * fil_vol; + + size_t limit = env::value(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK); + if(limit == 0) + limit = NAIVE_CONV_MAX_WORK; + + if(work > limit) + { + MIOPEN_LOG_I2("ConvDirectNaiveConv over work limit: work " + << work << " MACs exceeds limit " << limit + << " -- defer from Find benchmark when an alternative applies (avoids " + "TDR-prone naive launch)"); + return true; + } + return false; +} + /// Figure out the index of C (channel) stride so we can expand it into /// (G, C_per_group). Return value G_stride_idx is the position of G stride /// in the stride vector, such that the (G_stride_idx - 1) is the index that diff --git a/projects/miopen/test/gtest/unit_conv_solver_naive_work_gate.cpp b/projects/miopen/test/gtest/unit_conv_solver_naive_work_gate.cpp new file mode 100644 index 000000000000..d96b7af98ed0 --- /dev/null +++ b/projects/miopen/test/gtest/unit_conv_solver_naive_work_gate.cpp @@ -0,0 +1,199 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2026 Advanced Micro Devices, Inc. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + * + *******************************************************************************/ + +// Pure-CPU unit tests for the naive-conv work-size threshold +// (ConvDirectNaiveConvExceedsWorkLimit). Naive stays *applicable* at any size +// (it is the universal fallback); this threshold is consumed at the selection +// layer (EvaluateInvokers, src/conv/solver_finders.cpp) to keep the un-tiled +// naive kernel from being *benchmarked* on shapes so large its single launch +// would trip the OS GPU watchdog (TDR) when a non-naive alternative also +// applies. See src/solver/conv/conv_direct_naive_conv.cpp for the metric. +// +// These tests exercise the metric directly on ProblemDescriptions. They never +// construct a Handle and never launch a kernel, so they are safe on machines +// where the huge shapes would otherwise TDR. + +#include + +#include + +#include +#include + +#include "unit_TensorDescriptor.hpp" +#include "unit_conv_ConvolutionDescriptor.hpp" +#include "lib_env_var.hpp" +#include "gtest_common.hpp" + +namespace { + +using miopen::unit_tests::ConvolutionDescriptorParams; +using miopen::unit_tests::TensorDescriptorParams; +using Direction = miopen::conv::Direction; + +// Mirror of the env override read by the gate (declared UINT64 in the library). +MIOPEN_LIB_ENV_VAR(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK) + +// Build a 2D conv ProblemDescription with square spatial dims, odd square +// filter, unit stride/dilation and 'same' padding (so out spatial == in +// spatial). Layout/dtype are irrelevant to the work metric, so we use the +// simplest (NCHW, fp32); lens are always in NCHW logical order. +miopen::conv::ProblemDescription Make2dProblem( + std::size_t n, std::size_t c, std::size_t hw, std::size_t k, std::size_t fyx, Direction dir) +{ + const int pad = static_cast(fyx / 2); + TensorDescriptorParams in{miopenFloat, {n, c, hw, hw}}; + TensorDescriptorParams wei{miopenFloat, {k, c, fyx, fyx}}; + TensorDescriptorParams out{miopenFloat, {n, k, hw, hw}}; + ConvolutionDescriptorParams conv{{pad, pad}, {1, 1}, {1, 1}}; + return miopen::conv::ProblemDescription{in.GetTensorDescriptor(), + wei.GetTensorDescriptor(), + out.GetTensorDescriptor(), + conv.GetConvolutionDescriptor(), + dir}; +} + +// 3D variant (NCDHW), same conventions. +miopen::conv::ProblemDescription Make3dProblem( + std::size_t n, std::size_t c, std::size_t dhw, std::size_t k, std::size_t fzyx, Direction dir) +{ + const int pad = static_cast(fzyx / 2); + TensorDescriptorParams in{miopenFloat, {n, c, dhw, dhw, dhw}}; + TensorDescriptorParams wei{miopenFloat, {k, c, fzyx, fzyx, fzyx}}; + TensorDescriptorParams out{miopenFloat, {n, k, dhw, dhw, dhw}}; + ConvolutionDescriptorParams conv{{pad, pad, pad}, {1, 1, 1}, {1, 1, 1}}; + return miopen::conv::ProblemDescription{in.GetTensorDescriptor(), + wei.GetTensorDescriptor(), + out.GetTensorDescriptor(), + conv.GetConvolutionDescriptor(), + dir}; +} + +// Depthwise 2D conv (group_count == in_channels == out_channels), so +// C_per_group == 1. Same spatial conventions as Make2dProblem; weights are +// {k, C_per_group=1, fyx, fyx}. Because the work metric folds in the group count +// (below), every real depthwise conv lands far under the limit, so the selection +// gate never work-defers it -- Naive keeps competing on merit with no special +// case. These tests lock that safety invariant in. +miopen::conv::ProblemDescription +MakeDepthwise2dProblem(std::size_t n, std::size_t c, std::size_t hw, std::size_t fyx, Direction dir) +{ + const int pad = static_cast(fyx / 2); + const int group = static_cast(c); // depthwise: group == c == k + TensorDescriptorParams in{miopenFloat, {n, c, hw, hw}}; + TensorDescriptorParams wei{miopenFloat, {c, 1, fyx, fyx}}; // {k, C_per_group, fyx, fyx} + TensorDescriptorParams out{miopenFloat, {n, c, hw, hw}}; + ConvolutionDescriptorParams conv{{pad, pad}, {1, 1}, {1, 1}, group}; + return miopen::conv::ProblemDescription{in.GetTensorDescriptor(), + wei.GetTensorDescriptor(), + out.GetTensorDescriptor(), + conv.GetConvolutionDescriptor(), + dir}; +} + +bool ExceedsWorkLimit(const miopen::conv::ProblemDescription& p) +{ + return miopen::solver::conv::ConvDirectNaiveConvExceedsWorkLimit(p); +} + +} // namespace + +// The confirmed-TDR SDXL VAE decode conv (c128 k128 768^2 3x3, N=1) is ~87 GMAC, +// far above the ~16 GMAC default limit -> over threshold (naive skipped in the +// mini-bench when a non-naive alternative also applies). +TEST(CPU_ConvNaiveWorkGate_NONE, HugeShapeExceedsLimit) +{ + EXPECT_TRUE(ExceedsWorkLimit(Make2dProblem(1, 128, 768, 128, 3, Direction::Forward))); +} + +// A resnet50-class 3x3 (~0.1 GMAC) is orders of magnitude below the limit -> under +// threshold (naive runs normally in the mini-bench). +TEST(CPU_ConvNaiveWorkGate_NONE, SmallShapeWithinLimit) +{ + EXPECT_FALSE(ExceedsWorkLimit(Make2dProblem(1, 64, 56, 64, 3, Direction::Forward))); +} + +// A large 3D conv (~98 GMAC) is over threshold too -- the metric folds in output +// depth and the 3D filter volume. +TEST(CPU_ConvNaiveWorkGate_NONE, HugeShape3dExceedsLimit) +{ + EXPECT_TRUE(ExceedsWorkLimit(Make3dProblem(1, 64, 96, 64, 3, Direction::Forward))); +} + +// The MAC total is identical for fwd/bwd/wrw, so one constant governs all three +// directions: the huge shape is over threshold and the small shape is not, +// regardless of direction. +TEST(CPU_ConvNaiveWorkGate_NONE, DirectionInvariant) +{ + for(const auto dir : {Direction::Forward, Direction::BackwardData, Direction::BackwardWeights}) + { + EXPECT_TRUE(ExceedsWorkLimit(Make2dProblem(1, 128, 768, 128, 3, dir))); + EXPECT_FALSE(ExceedsWorkLimit(Make2dProblem(1, 64, 56, 64, 3, dir))); + } +} + +// Depthwise safety invariant for the selection gate: because the work metric folds +// in the group count (C_per_group == 1 for depthwise), the same c/k/hw/filter that +// is far *over* the limit as a dense conv is ~C times smaller as depthwise and lands +// well *under* it. This is why Naive is never work-deferred for depthwise and keeps +// competing (TDR-safe). A large-spatial depthwise (c=k=g=1024, 64^2, 3x3 ~= 38 MMAC) +// stays under threshold; the matching dense conv (~38 GMAC) is over. +TEST(CPU_ConvNaiveWorkGate_NONE, DepthwiseFarUnderLimit) +{ + EXPECT_FALSE(ExceedsWorkLimit(MakeDepthwise2dProblem(1, 1024, 64, 3, Direction::Forward))); +} + +TEST(CPU_ConvNaiveWorkGate_NONE, DenseCounterpartOverLimit) +{ + // Identical c/k/hw/filter as the depthwise above but dense (group == 1): ~C + // times more work, pushing it over the limit -- confirms the metric is + // group-aware and the depthwise case is not merely small by coincidence. + EXPECT_TRUE(ExceedsWorkLimit(Make2dProblem(1, 1024, 64, 1024, 3, Direction::Forward))); +} + +// The env override raises the limit; a huge shape that is normally over threshold +// falls under it when the cap is set above its work count (how a user with a raised +// TdrDelay opts naive back into the mini-bench on big shapes). +TEST(CPU_ConvNaiveWorkGate_NONE, EnvOverrideRaisesLimit) +{ + const auto huge = Make2dProblem(1, 128, 768, 128, 3, Direction::Forward); + ASSERT_TRUE(ExceedsWorkLimit(huge)); // over threshold with the default limit + ScopedEnvironment raise(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK, + std::uint64_t{1} << 60); + EXPECT_FALSE(ExceedsWorkLimit(huge)); +} + +// The env override also lowers the limit: a small shape that is normally under +// threshold goes over it when the cap is dropped below its work count. Confirms the +// override is honored in both directions rather than only disabling the check. +TEST(CPU_ConvNaiveWorkGate_NONE, EnvOverrideLowersLimit) +{ + const auto small = Make2dProblem(1, 64, 56, 64, 3, Direction::Forward); + ASSERT_FALSE(ExceedsWorkLimit(small)); // under threshold with the default limit + ScopedEnvironment lower(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK, + std::uint64_t{1000}); + EXPECT_TRUE(ExceedsWorkLimit(small)); +} From 7f7a941b68e7928569da4cef114d8c056d41e05d Mon Sep 17 00:00:00 2001 From: Brad Pepers Date: Sun, 23 Aug 2026 03:07:26 -0600 Subject: [PATCH 2/4] fix(miopen): add naive fallback when all non-naive solvers fail eval The naive-conv work gate skips *benchmarking* the un-tiled naive solver during Find when a non-naive solver is applicable and the problem exceeds the work limit. Applicability is not success, however: every applicable non-naive candidate can still be rejected at evaluation (e.g. insufficient workspace). Previously that left FindCore with an empty result and the convolution failed outright ("No suitable algorithm"), even though naive could have served it -- weakening the intended universal-fallback guarantee. FindCore now re-evaluates with the skip disabled when the result would otherwise be empty, so naive remains the true universal fallback. This only triggers in the rare empty-result case; the already-rejected non-naive candidates are re-rejected cheaply (workspace-filtered before any launch), and the slow naive fallback is left non-optimal so it is not persisted to find-db. The TDR win is preserved: naive is still skipped whenever any solver actually succeeds. Also: - Reword the retention log from the misleading "Unable to Skip Naive Solver" to a neutral "Retaining Naive Solver ()". - Extract IsNaiveSolverId to dedupe the naive-id string test. - CHANGELOG: record removal of MIOPEN_CONV_DIRECT_MAX_SIZE. Co-Authored-By: Claude Opus 4.8 --- projects/miopen/src/conv/solver_finders.cpp | 31 +++++++++++++++++++++ 1 file changed, 31 insertions(+) diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index fd8b5538c227..cf0bd7e91364 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -759,6 +759,37 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, std::make_move_iterator(evaluated.end())); } + // Universal-fallback guarantee. The work gate skips *benchmarking* Naive when a + // non-Naive solver is applicable, but applicability is not success: every non-Naive + // candidate can still be rejected at evaluation (e.g. insufficient workspace). If + // that leaves no solution at all, we would otherwise fail the convolution outright + // even though the un-tiled Naive kernel could have served it. Re-evaluate with the + // skip disabled so Naive remains the true universal fallback. This only triggers in + // the rare empty-result case, and the already-rejected non-Naive candidates are + // re-rejected cheaply (workspace-filtered before any kernel launch). + if(ret.solutions.empty() && non_naive_exists && naive_exceeds_work) + { + MIOPEN_LOG_I("No solver survived evaluation; re-running with the Naive benchmark " + "re-enabled as a last-resort fallback."); + for(const auto& ss : solutions) + { + auto evaluated = EvaluateInvokers(handle, + ss.second, + ss.first, + network_config, + invoke_ctx, + ret, + force_attach_binary, + non_naive_succeeded, + /*non_naive_exists=*/false, + /*naive_exceeds_work=*/false); + + ret.solutions.insert(ret.solutions.end(), + std::make_move_iterator(evaluated.begin()), + std::make_move_iterator(evaluated.end())); + } + } + return ret; } From 769afc1e8a01c29ac2cf341c3eb43a85e462c1f4 Mon Sep 17 00:00:00 2001 From: Brad Pepers Date: Wed, 30 Sep 2026 16:23:49 -0600 Subject: [PATCH 3/4] refactor(miopen): unify slow-solver handling behind a single speed class MIOpen had four partly-overlapping mechanisms for "this solver is applicable but is a catastrophically bad choice here": MIOPEN_CONV_DIRECT_MAX_SIZE, MIOPEN_NAIVE_TIMEOUT, the MIOPEN_SEARCH_CUTOFF slow branch, and GetWti. Each acted too late, too bluntly, or without the cross-solver view the decision needs. Replace all four with one declaration point and one consumption rule. Declare: SolverInterface::GetSpeedClass(ctx, problem) returns Normal, Slow, or ExceedsLaunchBudget. Consume: Find's mini-benchmark and immediate mode's fallback both benchmark everything in the lowest class with an applicable member and defer the rest. Slow and ExceedsLaunchBudget are separate classes rather than one bool. With a single flag, "all candidates are slow" means nothing is skipped and naive is launched anyway -- and all-slow is most likely exactly when the problem is huge, i.e. precisely the regime the gate exists for. This is a pre-launch gate. Bounding how long Find *waits* cannot avert a TDR; the dispatch has already happened. A sole-candidate ExceedsLaunchBudget solver is still benchmarked, so no problem loses coverage. The naive threshold is derived from device capability (CUs, wavefront width, clock) through a measured MAC-per-lane-cycle rate, calibrated over 109 timed shape/direction points on gfx1151 (RDNA3.5, wave32) and gfx950 (CDNA4, wave64). Across the 39 pairs measured on both, gfx950 ran a median 5.04x faster against a 5.30x lane-cycle ratio, so the scaling holds to within ~5%. Removed: MIOPEN_CONV_DIRECT_MAX_SIZE; MIOPEN_NAIVE_TIMEOUT and its separate-stream machinery (stream_tracker, naive_timeout); the SEARCH_CUTOFF slow branch; a dead non_slow_succeeded out-parameter. Escape hatch: MIOPEN_DEBUG_DEFER_SLOW_SOLVERS=0 restores prior behaviour. Relates to #3509, #1860, #9513. Co-Authored-By: Claude Opus 5 --- projects/miopen/CHANGELOG.md | 8 +- .../miopen/docs/reference/env_variables.rst | 66 +-- projects/miopen/src/conv/solver_finders.cpp | 469 ++++++------------ projects/miopen/src/env_debug.cpp | 9 +- projects/miopen/src/hip/convolutionocl.cpp | 63 ++- projects/miopen/src/hip/handlehip.cpp | 101 ---- .../miopen/src/include/miopen/any_solver.hpp | 28 +- .../include/miopen/conv/solver_finders.hpp | 38 +- .../src/include/miopen/conv/solvers.hpp | 47 +- .../src/include/miopen/find_solution.hpp | 17 +- projects/miopen/src/include/miopen/handle.hpp | 15 - projects/miopen/src/include/miopen/solver.hpp | 44 +- .../miopen/solver/conv_direct_naive_conv.hpp | 33 +- .../src/include/miopen/stream_tracker.hpp | 74 --- .../solver/conv/conv_direct_naive_conv.cpp | 160 ++++-- .../conv/conv_direct_naive_conv_bwd.cpp | 7 + .../conv/conv_direct_naive_conv_fwd.cpp | 7 + .../conv/conv_direct_naive_conv_wrw.cpp | 7 + projects/miopen/src/solver/conv/gemm.cpp | 37 +- projects/miopen/src/solver/conv/gemm_bwd.cpp | 39 +- projects/miopen/src/solver/conv/gemm_wrw.cpp | 24 +- projects/miopen/test/gtest/naive_timeout.cpp | 213 -------- projects/miopen/test/gtest/stream_tracker.cpp | 304 ------------ .../miopen/test/gtest/test_categories.yaml | 4 - .../unit_conv_solver_naive_work_gate.cpp | 213 +++++++- 25 files changed, 830 insertions(+), 1197 deletions(-) delete mode 100644 projects/miopen/src/include/miopen/stream_tracker.hpp delete mode 100644 projects/miopen/test/gtest/naive_timeout.cpp delete mode 100644 projects/miopen/test/gtest/stream_tracker.cpp diff --git a/projects/miopen/CHANGELOG.md b/projects/miopen/CHANGELOG.md index 3e6b9d6dc594..6999e00456a8 100644 --- a/projects/miopen/CHANGELOG.md +++ b/projects/miopen/CHANGELOG.md @@ -8,12 +8,18 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj ### Added * [Conv] Added gfx950 depthwise backward-weights (fp16/bf16) and gfx1250 depthwise kernels to the `ConvHipConv` solver (hipconv v0.3.1). +* [Conv] Added the `MIOPEN_DEBUG_DEFER_SLOW_SOLVERS` environment variable (default on). Setting it to 0 disables speed-class based solver deferral entirely, so find and immediate mode consider every applicable solver as they did before that mechanism existed. It is an escape hatch for a solver being deferred on a shape where it was in fact the right choice, and re-exposes the naive-kernel watchdog reset (TDR) the mechanism prevents. + +### Changed +* [Conv] Solvers that are applicable but known to be far off the pace for a given problem now declare this through `SolverInterface::GetSpeedClass`, which returns `Normal`, `Slow` (wasteful to benchmark, but safe to launch) or `ExceedsLaunchBudget` (a single launch is estimated to outlast the OS GPU watchdog). Both selection paths honor it: find's mini-benchmark and immediate mode's fallback now evaluate only the best class that has an applicable member and defer the rest. Ranking is preserved within a class, so when several equally-safe candidates remain they are all benchmarked and the least slow wins; separating `Slow` from `ExceedsLaunchBudget` is what keeps a watchdog risk from being launched merely because the alternative was also off the pace. A solver is never dropped when it is the only candidate, so no problem loses coverage. The naive convolution solvers report `ExceedsLaunchBudget` via the new work gate below; the explicit-GEMM solvers report `Slow` via their existing heuristics, which previously took effect only under the opt-in `MIOPEN_SEARCH_CUTOFF`. ### Resolved Issues -* [Conv] Fixed a GPU driver watchdog reset (TDR) that could occur while find benchmarked the un-tiled naive convolution solver on very large problems (e.g. VAE-decode and 3D convolutions). Naive is now skipped during find when a non-naive solver is applicable and the problem's total MAC work exceeds a threshold (~16 GMAC, overridable via `MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK`). This complements the existing `MIOPEN_NAIVE_TIMEOUT` deferral, which bounds how long find waits on the naive warmup but cannot prevent the launch itself. Naive still runs when it is the only applicable solver, and small problems where it is fastest continue to compete. +* [Conv] Fixed a GPU driver watchdog reset (TDR) that could occur while find benchmarked the un-tiled naive convolution solver on very large problems (e.g. VAE-decode and 3D convolutions). The three naive solvers now report `SolverSpeedClass::ExceedsLaunchBudget` when the problem's total MAC work exceeds what the device can retire inside a launch-time budget, so they are skipped during find and immediate mode whenever any better-classed solver applies. The limit is derived from device compute units, wavefront width, and clock rather than a fixed constant, so it scales across the CDNA/RDNA range; override the budget with `MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS` (default 1000 ms) or pin the MAC limit outright with `MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK`. Naive still runs when it is the only applicable solver, and small problems where it is fastest continue to compete. ### Removed * [Conv] Removed the `MIOPEN_CONV_DIRECT_MAX_SIZE` environment variable (an output-element-count threshold that disabled the entire Direct algorithm). It was a poor proxy for kernel launch cost and is superseded by the naive-conv work gate described above. +* [Conv] Removed the `MIOPEN_NAIVE_TIMEOUT` and `MIOPEN_NAIVE_TIMEOUT_FACTOR` environment variables along with the separate-stream evaluation machinery behind them. That deferral measured the naive warmup only after dispatching it, so it could not prevent the launch that causes a TDR; the work gate above declines the launch up front and makes the timeout path unreachable. +* [Conv] Removed the slow-solver handling from `MIOPEN_SEARCH_CUTOFF`, which dropped a slow solver unconditionally and could therefore discard the only applicable candidate. That declaration is now `GetSpeedClass` and is consumed by the last-resort policy described above, which is on by default and never drops a sole candidate. `MIOPEN_SEARCH_CUTOFF` continues to bound measured tuning and sampling time. * [Conv] Removed gfx803 convolution solver `ConvBinWinogradRxSFused` and its kernel sources. * [Conv] Removed the gfx803 code paths and kernel sources from `ConvBinWinograd3x3U` and `ConvBinWinogradRxS`. * [Conv] Removed the gfx803 find-database files. diff --git a/projects/miopen/docs/reference/env_variables.rst b/projects/miopen/docs/reference/env_variables.rst index a699c7571c59..a3b7de42a6e3 100644 --- a/projects/miopen/docs/reference/env_variables.rst +++ b/projects/miopen/docs/reference/env_variables.rst @@ -104,39 +104,39 @@ and :doc:`Performance database <../conceptual/perfdb>`. - | 1: Enable | 0 or unset: Disable - * - | ``MIOPEN_NAIVE_TIMEOUT`` - | When enabled and a non-naive solver has already succeeded, - | naive solver evaluation is launched on a separate HIP stream - | with a wall-clock budget of ``MIOPEN_NAIVE_TIMEOUT_FACTOR`` - | percent of the best non-naive time. If the naive kernel does - | not complete within the budget, it is skipped and the stream - | is returned to an async pool for later reclamation. Naive - | solvers are still used as a fallback when no non-naive solver - | succeeds. - | - | Limitation: a skipped kernel keeps reading the caller's input - | tensors until it finishes. Its results are discarded, so stale - | or reused data is harmless, but unmapping those pages while it - | reads them faults the process. Freeing with ``hipFree`` is safe - | because it synchronizes the device first. ``hipFreeAsync`` is - | ordered only against the stream it is given and does not wait - | for a skipped kernel; the pool normally keeps the address range - | reserved, but it releases pages back to the operating system - | once ``hipMemPoolAttrReleaseThreshold`` is exceeded, which its - | default value of 0 always is. Disable this option if the - | application frees convolution inputs through the - | stream-ordered allocator. - - | 1 or unset: Enable (default) - | 0: Disable - - * - | ``MIOPEN_NAIVE_TIMEOUT_FACTOR`` - | Sets the wall-clock budget for naive solver evaluation as a - | percentage of the best non-naive solver time. Only takes - | effect when ``MIOPEN_NAIVE_TIMEOUT`` is enabled and a - | non-naive solver has already succeeded. For example, the - | default value of 300 gives naive a budget of 3× the best - | non-naive time. - - | Integer percentage (default: 300) + * - | ``MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS`` + | Wall-clock budget, in milliseconds, for a single launch of the + | un-tiled naive convolution kernel. The naive solvers convert + | this budget into a MAC limit using the device's compute units, + | wavefront width, and clock, then report the speed class + | ``ExceedsLaunchBudget`` for any problem exceeding it, which + | keeps find and immediate mode from selecting them while any + | better-classed solver applies. Raise it to let + | naive compete on larger problems, at the risk of a GPU driver + | watchdog reset (TDR); lower it on a device with a short + | watchdog. Values above 60000 are clamped. + - | Milliseconds (default: 1000) + + * - | ``MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK`` + | Pins the naive convolution MAC limit to a fixed value instead + | of deriving it from the device and the time budget above. Takes + | precedence over ``MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS``. + | Intended for testing and for bisecting the derived limit. + - | MAC count (unset: derive from device) + + * - | ``MIOPEN_DEBUG_DEFER_SLOW_SOLVERS`` + | Master switch for speed-class based solver deferral. When on, + | each solver reports a speed class for the problem at hand and + | find and immediate mode consider only the best class that has + | an applicable member, so a solver that is merely applicable is + | not benchmarked or offered while a better-classed one exists. + | Set to 0 to consider every applicable solver regardless of its + | speed class, as releases before this mechanism did. That is a + | workaround for a solver being deferred on a shape where it was + | in fact the right choice; it also re-exposes the naive-kernel + | GPU driver watchdog reset (TDR) the mechanism prevents. + - | 1 or unset: Defer slow solvers + | 0: Consider all applicable solvers * - | ``MIOPEN_DEBUG_DISABLE_FIND_DB`` | Disables FindDb functionality. diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index cf0bd7e91364..fc3d5c532354 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -4,17 +4,20 @@ #include #include -#include +#include #include -#include +#include +#include #include #include #include #include +#include #include #include +#include #include #include #include @@ -23,7 +26,6 @@ #include #include #include -#include #include MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_GEMM) @@ -32,15 +34,20 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_WINOGRAD) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_FFT) MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_DEVICE_ARCH) -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_COMPILE_ONLY) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_FIND_CONV_INSUFFICIENT_WORKSPACE_ALLOW_FINDDB_UPDATE) -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_NAIVE_TIMEOUT, true) -MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_NAIVE_TIMEOUT_FACTOR, 300) -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_SEARCH_CUTOFF, false) +// MIOPEN_DEBUG_COMPILE_ONLY and MIOPEN_SEARCH_CUTOFF come from generic_search_controls.hpp. MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_FIND_SKIP_PCT, 130) +/// Master kill switch for the whole speed-class mechanism. Default on; set to 0 to restore +/// the pre-gate behaviour of benchmarking and offering every applicable solver regardless +/// of how slow it declares itself to be. Present so that a mis-classification -- a solver +/// deferred on a shape where it was in fact the right choice -- is a one-variable +/// workaround rather than a rebuild, at the cost of re-exposing the watchdog reset (TDR) +/// this mechanism exists to prevent. +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_DEFER_SLOW_SOLVERS) + namespace miopen { namespace conv { @@ -212,226 +219,88 @@ const std::vector>& GetConvSolverFinders() } // namespace conv -namespace { - -struct NaiveWarmup -{ - enum class Status - { - Completed, - TimedOut, - ScratchUnavailable, - UnsupportedInvokeParams, - }; - - Status status; - float elapsed; -}; - -/// Redirects the handle onto a tracker-owned stream with profiling off, restoring -/// both on scope exit so no early return can strand the handle off the root stream. -struct AutoExclusiveStream -{ - AutoExclusiveStream(const Handle& h, hipStream_t stream) - : handle(h), prev_profiling(h.IsProfilingEnabled()) - { - handle.SetExclusiveStream(stream); - handle.EnableProfiling(false); - } - - ~AutoExclusiveStream() - { - handle.SetExclusiveStream(nullptr); - handle.EnableProfiling(prev_profiling); - } - - AutoExclusiveStream(const AutoExclusiveStream&) = delete; - AutoExclusiveStream& operator=(const AutoExclusiveStream&) = delete; - -private: - const Handle& handle; - bool prev_profiling; -}; - -std::string NaiveSkipReason(NaiveWarmup::Status status, float best_time) -{ - switch(status) - { - case NaiveWarmup::Status::TimedOut: - return "exceeded " + std::to_string(env::value(MIOPEN_NAIVE_TIMEOUT_FACTOR)) + - "% of best non-naive time (" + std::to_string(best_time) + " ms)"; - case NaiveWarmup::Status::ScratchUnavailable: - return "output tensor exceeds the scratch buffer cap"; - case NaiveWarmup::Status::UnsupportedInvokeParams: - return "invoke params type does not support scratch redirection"; - case NaiveWarmup::Status::Completed: break; - } - return "completed"; -} - -} // namespace - -static NaiveWarmup TryNaiveWithTimeout(const Handle& handle, - const Invoker& invoker, - const AnyInvokeParams& invoke_ctx, - float best_time) -{ - std::shared_ptr scratch; - AnyInvokeParams scratch_ctx; - - if(invoke_ctx.IsOfType()) - { - auto params = invoke_ctx.CastTo(); - scratch = handle.GetScratchBuffer(params.tensors.outDesc.GetNumBytes()); - if(!scratch) - return {NaiveWarmup::Status::ScratchUnavailable, 0.0f}; - params.tensors.out = scratch->buffer.get(); - scratch_ctx = AnyInvokeParams{params}; - } - else if(invoke_ctx.IsOfType()) - { - auto params = invoke_ctx.CastTo(); - scratch = handle.GetScratchBuffer(params.tensors.dwDesc.GetNumBytes()); - if(!scratch) - return {NaiveWarmup::Status::ScratchUnavailable, 0.0f}; - params.tensors.dw = scratch->buffer.get(); - scratch_ctx = AnyInvokeParams{params}; - } - else - return {NaiveWarmup::Status::UnsupportedInvokeParams, 0.0f}; - - auto& tracker = handle.GetStreamTracker(); - auto slot = tracker.acquire(handle); - slot.scratch = scratch; - - AutoExclusiveStream stream_guard{handle, slot.stream}; - - HipEventPtr ev_start = make_hip_event(); - HipEventPtr ev_stop = make_hip_event(); - - try - { - auto ev_status = hipEventRecord(ev_start.get(), slot.stream); - if(ev_status != hipSuccess) - MIOPEN_THROW_HIP_STATUS(ev_status, "Failed to record naive start event"); - invoker(handle, scratch_ctx); - ev_status = hipEventRecord(ev_stop.get(), slot.stream); - if(ev_status != hipSuccess) - MIOPEN_THROW_HIP_STATUS(ev_status, "Failed to record naive stop event"); - } - catch(...) - { - tracker.abandon(slot); - throw; - } - - const float timeout_factor = - static_cast(env::value(MIOPEN_NAIVE_TIMEOUT_FACTOR)) / 100.0f; - const float naive_budget = best_time * timeout_factor; - const auto deadline = std::chrono::steady_clock::now() + - std::chrono::microseconds(static_cast(naive_budget * 1000)); - bool finished = false; - while(std::chrono::steady_clock::now() < deadline) - { - if(hipEventQuery(ev_stop.get()) == hipSuccess) - { - finished = true; - break; - } - std::this_thread::yield(); - } +bool IsDeferSlowSolversEnabled() { return !env::disabled(MIOPEN_DEBUG_DEFER_SLOW_SOLVERS); } - if(finished) - { - tracker.release(slot); - float warmup_elapsed = 0.0f; - (void)hipEventElapsedTime(&warmup_elapsed, ev_start.get(), ev_stop.get()); - return {NaiveWarmup::Status::Completed, warmup_elapsed}; - } - - tracker.abandon(slot); - return {NaiveWarmup::Status::TimedOut, 0.0f}; -} - -/// A solver is Naive iff its id contains "Naive" (ConvDirectNaiveConv{Fwd,Bwd,Wrw}). -/// Single source of truth for the string test used by the per-solver skip check, the -/// timeout-deferral ordering, and the FindCore "does any non-Naive solver apply" scan. -static bool IsNaiveSolverId(const std::string& solver_id) +/// The speed class a solver reported for this problem, defaulting to Normal for anything +/// the caller did not record (which is everything, when the policy is disabled). +static solver::SolverSpeedClass +SpeedClassOf(const std::map& speed_classes, + const std::string& solver_id) { - return solver_id.find("Naive") != std::string::npos; + const auto it = speed_classes.find(solver_id); + return it == speed_classes.end() ? solver::SolverSpeedClass::Normal : it->second; } -/// Decide whether the Naive convolution solver should be *skipped* -- i.e. not -/// executed in Find's mini-benchmark -- for the current solution. +/// Decide whether a solver should be skipped -- i.e. not executed at all in Find's +/// mini-benchmark -- because something better-classed is available. +/// +/// Why this exists: some solvers are deliberately *applicable* far outside the region +/// where they are a sensible choice. Naive convolution is MIOpen's universal fallback and +/// must stay applicable at any problem size, but its kernel is un-tiled, so on a large +/// problem a single launch runs for multiple seconds. Explicit GEMM stays applicable into +/// shapes where im2col expansion and per-batch BLAS calls make it wildly uncompetitive. +/// Find times every applicable solver by actually *executing* it, so merely benchmarking +/// such a solver costs real time -- and in the naive case trips the OS GPU watchdog +/// (a TDR / driver reset) even when a fast solver also applies and would ultimately win. +/// The solver does not have to be *selected* to cause the hang; being *benchmarked* is +/// enough. /// -/// Why this exists: Naive is MIOpen's universal fallback and stays *applicable* at -/// any problem size, but its kernel is un-tiled, so on a large problem a single -/// launch runs for multiple seconds. Find times every applicable solver by actually -/// *executing* it, so merely benchmarking Naive on such a shape trips the OS GPU -/// watchdog (a TDR / driver reset) even when a fast solver (CK/GEMM) also applies -/// and would ultimately win. Naive does not have to be *selected* to cause the -/// hang -- being *benchmarked* is enough. +/// A solver classifies itself via SolverInterface::GetSpeedClass(ctx, problem), which sees +/// the concrete problem and the device. That keeps the *criterion* with the solver that +/// understands it (MAC work vs. a device-derived limit for naive, arch-aware shape +/// heuristics for GEMM) while this function owns the one thing a per-solver hook cannot +/// see: what else is available. /// -/// This is deliberately a *pre-launch* gate, complementing (not replacing) the -/// TryNaiveWithTimeout path below. That path bounds how long Find *waits* on the -/// naive warmup, but the dispatch has already happened and the kernel keeps -/// occupying the GPU after the slot is abandoned -- which is exactly what the OS -/// watchdog measures, so abandoning the wait does not avert a TDR. It also only -/// engages once a non-Naive solver has actually *succeeded* (non_naive_succeeded), -/// so it does nothing when every alternative is rejected at evaluation -- the leak -/// this gate closes by keying off applicability instead. We therefore skip -/// *launching* Naive during the benchmark exactly when BOTH of these hold: +/// The rule is a single comparison: benchmark everything in the best (lowest) class that +/// has an applicable member, skip everything worse. benchmark_class is that class, +/// computed by the caller over the full solution list across *all* algorithms and *before* +/// workspace filtering -- so a fast-but-workspace-limited alternative still counts, which +/// is what closes the leak where the alternative never gets to time itself. /// -/// * non_naive_exists -- some non-Naive solver was found applicable for this -/// problem. The caller computes this from the full -/// solution list *before* workspace filtering, so a fast -/// solver that is later workspace-filtered still counts; -/// that is what closes the TDR leak where the alternative -/// never gets to time itself. AND -/// * naive_exceeds_work -- the problem's total MAC work is over the Naive work -/// limit (~16 GMAC, see ConvDirectNaiveConvExceedsWorkLimit), -/// i.e. large enough that a single launch could TDR. +/// Consequences of that one rule (the whole policy in one place): +/// * Any Normal candidate exists -> every deferred solver is skipped and behaviour is +/// exactly as if this mechanism did not exist. +/// * No Normal, several Slow -> all the Slow ones are benchmarked and the least slow +/// wins on measurement. Deferral ranks candidates; it does not discard a whole tier +/// in favour of an arbitrary survivor. +/// * Only ExceedsLaunchBudget left -> it is benchmarked, preserving the universal- +/// fallback guarantee. A huge sole-naive shape may then still TDR -- an honest +/// "extend coverage here" signal we deliberately do not mask. +/// * Crucially, Slow and ExceedsLaunchBudget are *separate* classes rather than one +/// "slow" flag, so a merely wasteful solver outranks a watchdog risk: naive is never +/// launched just because GEMM also declared itself off the pace. With a single flag +/// that case -- all candidates slow, which is most likely exactly when the problem is +/// huge -- would have fallen through the gate and launched naive anyway. +/// * A solver outside its deferred regime (Normal) competes and wins on merit wherever +/// it is fastest. For naive this notably covers *all* real depthwise convs +/// (group == C == K => c_per_group == 1 => ~C x less work), which never approach the +/// work limit and for which naive is often the fastest option: they keep competing +/// with no special case. /// -/// Consequences of this exact condition (the whole selection policy in one place): -/// * Naive is the *sole* applicable solver -> not skipped -> it runs, preserving -/// the universal-fallback guarantee. A huge sole-Naive shape may then still TDR -/// -- an honest "extend coverage here" signal we deliberately do not mask. -/// * Any shape under the work limit -> not skipped -> Naive competes and wins on -/// merit wherever it is fastest. This notably covers *all* real depthwise convs -/// (group == C == K => C_per_group == 1 => ~C x less work), which never approach -/// the limit and for which Naive is often the fastest option (e.g. NCHW): they -/// keep competing with no special case. -/// * Only large-and-avoidable shapes (over the limit with an alternative present) -/// are skipped -- the actual TDR case. +/// This is deliberately a *pre-launch* gate. Bounding how long Find *waits* on a slow +/// solver cannot avert a TDR, because the dispatch has already happened and the kernel +/// keeps occupying the GPU after the wait is abandoned -- which is exactly what the OS +/// watchdog measures. /// -/// MIOPEN_NAIVE_TIMEOUT (naive_timeout, default on) is the master switch for Naive -/// deferral and gates this policy as well as the timeout path; unset it to force -/// Naive to always compete. The work limit itself is separately tunable via -/// MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK (set it very large to disable just this -/// gate while keeping the timeout). Golden references are unaffected: GpuConvReference -/// compiles and launches the naive kernel directly, bypassing the solver framework, so -/// verification never reaches this gate. -static bool ShouldSkipNaiveBenchmark(bool is_naive, - bool naive_timeout, - bool non_naive_exists, - bool naive_exceeds_work) +/// Golden references are unaffected: GpuConvReference compiles and launches the naive +/// kernel directly, bypassing the solver framework, so verification never reaches here. +static bool ShouldSkipSlowBenchmark(solver::SolverSpeedClass speed_class, + solver::SolverSpeedClass benchmark_class) { - if(!is_naive || !naive_timeout) - return false; - return non_naive_exists && naive_exceeds_work; + return speed_class > benchmark_class; } /// Register invoker only for the best solution within algorithm. -std::vector EvaluateInvokers(const Handle& handle, - const std::vector& solutions, - const AlgorithmName& algorithm_name, - const NetworkConfig& network_config, - const AnyInvokeParams& invoke_ctx, - FindCoreResult& core_result, - bool force_attach_binary, - bool& non_naive_succeeded, - bool non_naive_exists, - bool naive_exceeds_work) +std::vector +EvaluateInvokers(const Handle& handle, + const std::vector& solutions, + const AlgorithmName& algorithm_name, + const NetworkConfig& network_config, + const AnyInvokeParams& invoke_ctx, + FindCoreResult& core_result, + bool force_attach_binary, + const std::map& speed_classes, + solver::SolverSpeedClass benchmark_class) { std::vector ret; @@ -439,56 +308,44 @@ std::vector EvaluateInvokers(const Handle& handle, if(!arch.empty()) return ret; - const auto is_naive_solver = [](const solver::ConvSolution& s) { - return IsNaiveSolverId(s.solver_id); + const auto speed_class_of = [&](const solver::ConvSolution& s) { + return SpeedClassOf(speed_classes, s.solver_id); }; - bool naive_timeout = env::value(MIOPEN_NAIVE_TIMEOUT); bool using_search_cutoff = env::value(MIOPEN_SEARCH_CUTOFF); - // Defer Naive only when a non-Naive alternative exists across all algorithms or this one. - const bool defer_naive = - naive_timeout && - (non_naive_succeeded || std::any_of(solutions.begin(), solutions.end(), [&](const auto& s) { - return !is_naive_solver(s); - })); + auto selected = miopen::solver::ConvSolution{miopenStatusUnknownError}; auto best = std::numeric_limits::max(); auto best_invoker = Invoker{}; std::vector samples; - // Iterate non-Naive solutions first, Naive last + // Benchmark in speed-class order, best first, so that the cheap candidates establish + // find_search_best_time (and hence the measured-time cutoff below) before an expensive + // one is timed against it. Stable, so solvers within a class keep finder order. std::vector order(solutions.size()); std::iota(order.begin(), order.end(), 0); - if(defer_naive) - { - std::stable_partition(order.begin(), order.end(), [&](std::size_t i) { - return !is_naive_solver(solutions[i]); - }); - } + std::stable_sort(order.begin(), order.end(), [&](std::size_t a, std::size_t b) { + return speed_class_of(solutions[a]) < speed_class_of(solutions[b]); + }); for(std::size_t idx : order) { const auto& sol = solutions[idx]; - const bool is_naive = is_naive_solver(sol); - if(ShouldSkipNaiveBenchmark(is_naive, naive_timeout, non_naive_exists, naive_exceeds_work)) + const auto speed_class = speed_class_of(sol); + if(ShouldSkipSlowBenchmark(speed_class, benchmark_class)) { - MIOPEN_LOG_I("Skipping Naive Solver: " << algorithm_name.ToString() << ":" - << sol.solver_id); + MIOPEN_LOG_I("Skipping last-resort (slow) Solver: " << algorithm_name.ToString() << ":" + << sol.solver_id); continue; } - if(naive_timeout && is_naive) + if(speed_class != solver::SolverSpeedClass::Normal) { - // Naive is being kept in the benchmark set even though the work gate is on. - // Name the reason so this reads as an expected retention, not an anomaly. It may - // still be cut short by the timeout path below (cutoff_naive). - const auto* reason = !non_naive_exists ? "sole applicable solver" - : !naive_exceeds_work ? "below work limit" - : "skip criteria not met"; - MIOPEN_LOG_I("Retaining Naive Solver (" << reason << "): " << algorithm_name.ToString() - << ":" << sol.solver_id); + // A last-resort solver is being benchmarked anyway because nothing better-classed + // applies. Name the reason so this reads as an expected retention, not an anomaly. + MIOPEN_LOG_I("Retaining last-resort (slow) Solver (no better alternative applies): " + << algorithm_name.ToString() << ":" << sol.solver_id); } - const bool cutoff_naive = defer_naive && is_naive && non_naive_succeeded; if(!conv::IsEnoughWorkspace( "EvaluateInvokers", solver::Id{sol.solver_id}, sol.workspace_sz, &invoke_ctx)) @@ -557,22 +414,6 @@ std::vector EvaluateInvokers(const Handle& handle, int i = 0; samples.clear(); - if(cutoff_naive) - { - const auto warmup = TryNaiveWithTimeout( - handle, invoker, invoke_ctx, core_result.find_search_best_time); - if(warmup.status != NaiveWarmup::Status::Completed) - { - MIOPEN_LOG_I( - "Skipped naive solver " - << algorithm_name.ToString() << ":" << sol.solver_id << ": " - << NaiveSkipReason(warmup.status, core_result.find_search_best_time)); - continue; - } - first_elapsed = warmup.elapsed; - i = 1; - } - while(i < N_RUNS_MAX && elapsed < TIME_MS_MAX) { invoker(handle, invoke_ctx); @@ -639,8 +480,6 @@ std::vector EvaluateInvokers(const Handle& handle, else solution.SetInvoker(invoker, {}, {}); ret.emplace_back(std::move(solution)); - if(!is_naive) - non_naive_succeeded = true; } catch(const miopen::Exception& ex) { @@ -723,54 +562,62 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, ret.solutions.reserve(total); - // Does any non-Naive solver apply, across all algorithms? Computed from the full solution - // list *before* per-solver workspace filtering, so a fast-but-workspace-limited alternative - // still counts and Naive is deferred rather than benchmarked (the TDR-leak fix). Also decide - // once whether this conv's total MAC work is large enough that the un-tiled Naive kernel would - // trip the OS GPU watchdog if benchmarked. Non-conv (fusion) problems yield nullptr and are - // inert here. - const bool non_naive_exists = - std::any_of(solutions.begin(), solutions.end(), [](const auto& g) { - return std::any_of(g.second.begin(), g.second.end(), [](const auto& s) { - return !IsNaiveSolverId(s.solver_id); - }); - }); + // Ask every candidate how far off the pace it considers itself for *this* problem on + // *this* device (SolverInterface::GetSpeedClass), and record the ones that are not + // Normal. Computed here rather than inside EvaluateInvokers for two reasons: + // EvaluateInvokers sees one algorithm group at a time, whereas "what is the best class + // available" is a question about the whole candidate set; and this evaluates + // GetSpeedClass once per solver instead of once per benchmark iteration. + // + // The scan runs over the full solution list *before* per-solver workspace filtering, + // so a fast-but-workspace-limited alternative still counts and the deferred solver is + // skipped rather than benchmarked -- that is the leak this closes, since otherwise the + // alternative never gets to time itself and the deferred solver runs anyway. + // + // Non-conv (fusion) problems yield nullptr and leave the map empty, making the whole + // mechanism inert for them -- as does MIOPEN_DEBUG_DEFER_SLOW_SOLVERS=0. An empty map + // means every candidate reads back as Normal, so classes_present is {Normal}, the loop + // below runs exactly once and nothing is skipped: the pre-gate behaviour, restored by + // not gathering the data rather than by a second code path. + std::map speed_classes; const auto* conv_problem = dynamic_cast(&problem); - const bool naive_exceeds_work = - (conv_problem != nullptr) && - solver::conv::ConvDirectNaiveConvExceedsWorkLimit(*conv_problem); - - bool non_naive_succeeded = false; - for(const auto& ss : solutions) + if(conv_problem != nullptr && IsDeferSlowSolversEnabled()) { - auto evaluated = EvaluateInvokers(handle, - ss.second, - ss.first, - network_config, - invoke_ctx, - ret, - force_attach_binary, - non_naive_succeeded, - non_naive_exists, - naive_exceeds_work); - - ret.solutions.insert(ret.solutions.end(), - std::make_move_iterator(evaluated.begin()), - std::make_move_iterator(evaluated.end())); + for(const auto& group : solutions) + { + for(const auto& s : group.second) + { + const auto any = solver::Id{s.solver_id}.GetSolver(); + if(any.IsEmpty()) + continue; + const auto speed_class = any.GetSpeedClass(ctx, *conv_problem); + if(speed_class != solver::SolverSpeedClass::Normal) + speed_classes.emplace(s.solver_id, speed_class); + } + } } - // Universal-fallback guarantee. The work gate skips *benchmarking* Naive when a - // non-Naive solver is applicable, but applicability is not success: every non-Naive - // candidate can still be rejected at evaluation (e.g. insufficient workspace). If - // that leaves no solution at all, we would otherwise fail the convolution outright - // even though the un-tiled Naive kernel could have served it. Re-evaluate with the - // skip disabled so Naive remains the true universal fallback. This only triggers in - // the rare empty-result case, and the already-rejected non-Naive candidates are - // re-rejected cheaply (workspace-filtered before any kernel launch). - if(ret.solutions.empty() && non_naive_exists && naive_exceeds_work) + // The distinct classes present among the candidates, ascending (std::set is ordered), + // i.e. the sequence of benchmark_class values to try. Normally this is just {Normal} + // and the loop below runs once. + std::set classes_present; + for(const auto& group : solutions) + for(const auto& s : group.second) + classes_present.insert(SpeedClassOf(speed_classes, s.solver_id)); + if(classes_present.empty()) + classes_present.insert(solver::SolverSpeedClass::Normal); + + // Universal-fallback guarantee. The gate benchmarks only the best class present, but + // applicability is not success: every candidate in that class can still be rejected at + // evaluation (e.g. insufficient workspace). If that leaves no solution at all, we would + // otherwise fail the convolution outright even though a deferred solver could have + // served it. So descend one class at a time until something survives -- one class, not + // straight to "gate off", because dropping from Slow to ExceedsLaunchBudget wholesale + // is exactly the watchdog risk the classes exist to order. Past the first iteration + // this only triggers in the rare empty-result case, and the already-rejected candidates + // are re-rejected cheaply (workspace-filtered before any kernel launch). + for(const auto benchmark_class : classes_present) { - MIOPEN_LOG_I("No solver survived evaluation; re-running with the Naive benchmark " - "re-enabled as a last-resort fallback."); for(const auto& ss : solutions) { auto evaluated = EvaluateInvokers(handle, @@ -780,14 +627,20 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, invoke_ctx, ret, force_attach_binary, - non_naive_succeeded, - /*non_naive_exists=*/false, - /*naive_exceeds_work=*/false); + speed_classes, + benchmark_class); ret.solutions.insert(ret.solutions.end(), std::make_move_iterator(evaluated.begin()), std::make_move_iterator(evaluated.end())); } + + if(!ret.solutions.empty()) + break; + + if(benchmark_class != *classes_present.rbegin()) + MIOPEN_LOG_I("No solver survived evaluation; descending to the next (slower) " + "speed class as a fallback."); } return ret; diff --git a/projects/miopen/src/env_debug.cpp b/projects/miopen/src/env_debug.cpp index 3c5cca59e3a9..bbdd61b75b65 100644 --- a/projects/miopen/src/env_debug.cpp +++ b/projects/miopen/src/env_debug.cpp @@ -50,12 +50,14 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3) MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_DIRECT) +MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS) MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_FFT) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_GEMM) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_WINOGRAD) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_DEFER_SLOW_SOLVERS) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_DISABLE_SYSTEM_DB) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_DISABLE_USER_DB) MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_DEBUG_FIND_ONLY_SOLVER) @@ -74,8 +76,6 @@ MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_CUSTOM_CACHE_DIR) MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_LOG_LEVEL) MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_LOG_BUFFER_SIZE, 0); MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_DEBUG_CONV_DIRECT_ASM_WRW1X1_PERF_VALS) -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_NAIVE_TIMEOUT, true) -MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_NAIVE_TIMEOUT_FACTOR, 300) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_SKIP_ASAN_DISABLED_TESTS, true) namespace miopen { @@ -194,6 +194,8 @@ const LibEnvVar& FindEnvVariable(std::string_view name) {MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3.GetName(), MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3}, {MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS.GetName(), MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS}, {MIOPEN_DEBUG_CONV_DIRECT.GetName(), MIOPEN_DEBUG_CONV_DIRECT}, + {MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS.GetName(), + MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS}, {MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK.GetName(), MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK}, {MIOPEN_DEBUG_CONV_FFT.GetName(), MIOPEN_DEBUG_CONV_FFT}, @@ -202,6 +204,7 @@ const LibEnvVar& FindEnvVariable(std::string_view name) {MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS.GetName(), MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS}, {MIOPEN_DEBUG_CONV_WINOGRAD.GetName(), MIOPEN_DEBUG_CONV_WINOGRAD}, {MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL.GetName(), MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL}, + {MIOPEN_DEBUG_DEFER_SLOW_SOLVERS.GetName(), MIOPEN_DEBUG_DEFER_SLOW_SOLVERS}, {MIOPEN_DEBUG_DISABLE_SYSTEM_DB.GetName(), MIOPEN_DEBUG_DISABLE_SYSTEM_DB}, {MIOPEN_DEBUG_DISABLE_USER_DB.GetName(), MIOPEN_DEBUG_DISABLE_USER_DB}, {MIOPEN_DEBUG_FIND_ONLY_SOLVER.GetName(), MIOPEN_DEBUG_FIND_ONLY_SOLVER}, @@ -215,8 +218,6 @@ const LibEnvVar& FindEnvVariable(std::string_view name) {MIOPEN_FIND_ENFORCE.GetName(), MIOPEN_FIND_ENFORCE}, {MIOPEN_FIND_MODE.GetName(), MIOPEN_FIND_MODE}, {MIOPEN_FIND_MODE_FUSION.GetName(), MIOPEN_FIND_MODE_FUSION}, - {MIOPEN_NAIVE_TIMEOUT.GetName(), MIOPEN_NAIVE_TIMEOUT}, - {MIOPEN_NAIVE_TIMEOUT_FACTOR.GetName(), MIOPEN_NAIVE_TIMEOUT_FACTOR}, {MIOPEN_LOG_LEVEL.GetName(), MIOPEN_LOG_LEVEL}, {MIOPEN_LOG_BUFFER_SIZE.GetName(), MIOPEN_LOG_BUFFER_SIZE}, {MIOPEN_USER_DB_PATH.GetName(), MIOPEN_USER_DB_PATH}, diff --git a/projects/miopen/src/hip/convolutionocl.cpp b/projects/miopen/src/hip/convolutionocl.cpp index 444e625cf057..3f55f59f7121 100644 --- a/projects/miopen/src/hip/convolutionocl.cpp +++ b/projects/miopen/src/hip/convolutionocl.cpp @@ -29,6 +29,7 @@ #include #include +#include #include #include #include @@ -296,15 +297,10 @@ std::vector EvaluateConvSolutions(const ExecutionContext& ctx, AlgorithmName algo{ ConvolutionAlgoToDirectionalString(id.GetAlgo(), problem.GetDirection())}; - bool ocl_non_naive_succeeded = false; - std::vector eval_sol = EvaluateInvokers(handle, - conv_sols, - algo, - problem.MakeNetworkConfig(), - invoke_ctx, - core_result, - false, - ocl_non_naive_succeeded); + // Single-solver evaluation: the speed-class policy needs a cross-solver view it + // cannot have here, so it is left disabled (default empty class map). + std::vector eval_sol = EvaluateInvokers( + handle, conv_sols, algo, problem.MakeNetworkConfig(), invoke_ctx, core_result, false); if(!eval_sol.empty()) eval_sols.emplace_back(eval_sol.front()); @@ -1045,6 +1041,55 @@ ConvolutionDescriptor::GetSolutionsFallback(const ExecutionContext& ctx, interim.emplace_back(miopenConvSolution_t{wti2time(wti), ws, solver_id.Value(), algo}); } } + // Last-resort (slow) solver gate, immediate-mode half. Neither fallback path above + // measures anything -- the AI path assigns synthetic ranks and the WTI path a static + // estimate -- so nothing here would notice that a solver is applicable but wildly + // unsuitable for this shape until the caller runs it. For the un-tiled naive kernel + // that means seconds of GPU occupancy and an OS watchdog reset (a TDR); for explicit + // GEMM it means im2col expansion and per-batch BLAS calls nobody wanted. + // + // Apply the same rule Find applies in EvaluateInvokers: ask each candidate what + // SolverInterface::GetSpeedClass it reports for this problem and device, and keep only + // those in the best (lowest) class present. The cross-solver view has to live here + // rather than in the solver's own GetWti, because a per-solver hook cannot see whether + // an alternative exists and would therefore also drop the solver when it is the *only* + // candidate, turning a slow convolution into a failed one. Keeping a whole class, not + // just one winner, preserves ranking among equally-safe candidates; keeping only the + // best class is what stops a watchdog risk (ExceedsLaunchBudget) from being handed out + // merely because the alternative was also off the pace (Slow). Placed after both paths + // and before the sort so it covers whichever one populated interim. + // + // MIOPEN_DEBUG_DEFER_SLOW_SOLVERS=0 skips the whole block, leaving every applicable + // candidate on offer exactly as before the gate existed. + if(!interim.empty() && IsDeferSlowSolversEnabled()) + { + std::vector classes; + classes.reserve(interim.size()); + for(const auto& s : interim) + { + const auto any = solver::Id{s.solution_id}.GetSolver(); + classes.push_back(any.IsEmpty() ? solver::SolverSpeedClass::Normal + : any.GetSpeedClass(ctx, problem)); + } + + const auto best = *std::min_element(begin(classes), end(classes)); + if(best != *std::max_element(begin(classes), end(classes))) + { + const auto dropped = + std::count_if(begin(classes), end(classes), [&](auto c) { return c != best; }); + MIOPEN_LOG_I2("Immediate fallback: dropping " + << dropped << " last-resort (slow) candidate(s); " + << (interim.size() - dropped) << " better-classed candidate(s) remain"); + + auto kept = std::vector{}; + kept.reserve(interim.size() - dropped); + for(size_t i = 0; i < interim.size(); ++i) + if(classes[i] == best) + kept.push_back(interim[i]); + interim = std::move(kept); + } + } + MIOPEN_LOG_I2("maxSolutionCount = " << maxSolutionCount << ", available = " << interim.size()); for(const auto& s : interim) MIOPEN_LOG_I2(s); diff --git a/projects/miopen/src/hip/handlehip.cpp b/projects/miopen/src/hip/handlehip.cpp index db4f259addca..a4c6469ed5e7 100644 --- a/projects/miopen/src/hip/handlehip.cpp +++ b/projects/miopen/src/hip/handlehip.cpp @@ -172,9 +172,6 @@ int set_default_device() // NOLINTNEXTLINE (cppcoreguidelines-avoid-non-const-global-variables) static thread_local unsigned int miopenHandle_current_stream_id = 0; -// Overrides the stream-pool index when set. See Handle::SetExclusiveStream. -// NOLINTNEXTLINE (cppcoreguidelines-avoid-non-const-global-variables) -static thread_local hipStream_t miopenHandle_exclusive_stream = nullptr; struct HandleImpl { // typedef MIOPEN_MANAGE_PTR(hipStream_t, hipStreamDestroy) StreamPtr; @@ -286,9 +283,6 @@ struct HandleImpl Allocator allocator{}; KernelCache cache; TargetProperties target_properties; - mutable StreamTracker stream_tracker_; - std::weak_ptr active_scratch_; - std::size_t scratch_cap_ = 0; }; Handle::Handle(miopenAcceleratorQueue_t stream) : impl(std::make_unique()) @@ -397,21 +391,8 @@ void Handle::ReserveExtraStreamsInPool(int cnt) const } } -StreamTracker::StreamPtr Handle::CreateExclusiveStream() const -{ - impl->set_ctx(); - return impl->create_stream_non_blocking(); -} - -void Handle::SetExclusiveStream(miopenAcceleratorQueue_t stream) const -{ - miopenHandle_exclusive_stream = stream; -} - miopenAcceleratorQueue_t Handle::GetStream() const { - if(miopenHandle_exclusive_stream != nullptr) - return miopenHandle_exclusive_stream; if(miopenHandle_current_stream_id == 0) return impl->root_stream.get(); // locking only if handle in multistream mode @@ -419,79 +400,6 @@ miopenAcceleratorQueue_t Handle::GetStream() const return this->impl->ms_resourse_ptr->stream_pool.at(miopenHandle_current_stream_id - 1).get(); } -StreamTracker& Handle::GetStreamTracker() const { return impl->stream_tracker_; } - -std::shared_ptr Handle::GetScratchBuffer(std::size_t sz) const -{ - if(sz == 0) - return nullptr; - - if(impl->scratch_cap_ == 0) - impl->scratch_cap_ = GetGlobalMemorySize() / 8; - - if(sz > impl->scratch_cap_) - return nullptr; - - if(auto existing = impl->active_scratch_.lock(); existing && sz <= existing->size) - return existing; - - auto alloc = std::make_shared(); - alloc->buffer = impl->allocator(sz); - alloc->size = sz; - impl->active_scratch_ = alloc; - return alloc; -} - -void StreamTracker::sweep() -{ - if(draining_.empty()) - return; - - for(auto it = draining_.begin(); it != draining_.end();) - { - if(hipStreamQuery(it->stream) == hipSuccess) - { - it->scratch.reset(); - available_.push_back(std::move(*it)); - it = draining_.erase(it); - } - else - ++it; - } -} - -StreamTracker::~StreamTracker() -{ - if(draining_.empty()) - return; - - MIOPEN_LOG_I("Waiting for " << draining_.size() << " abandoned stream(s) to drain"); - - const auto start = std::chrono::steady_clock::now(); - for(auto& slot : draining_) - (void)hipStreamSynchronize(slot.stream); - const std::chrono::duration elapsed = - std::chrono::steady_clock::now() - start; - - MIOPEN_LOG_I("Drained " << draining_.size() << " abandoned stream(s) in " << elapsed.count() - << " ms"); -} - -StreamTracker::Slot StreamTracker::acquire(const Handle& handle) -{ - sweep(); - - if(!available_.empty()) - { - auto slot = std::move(available_.back()); - available_.pop_back(); - return slot; - } - - owned_streams_.push_back(handle.CreateExclusiveStream()); - return {owned_streams_.back().get(), {}}; -} - void Handle::SetAllocator(miopenAllocatorFunction allocator, miopenDeallocatorFunction deallocator, void* allocatorContext) const @@ -621,9 +529,6 @@ std::vector Handle::GetKernelsImpl(const std::string& algorithm, KernelInvoke Handle::Run(Kernel k, bool coop_launch) const { this->impl->set_ctx(); - // Reclaim scratch pinned by abandoned naive evaluations as soon as their - // streams go idle, rather than waiting for the next acquire(). - this->impl->stream_tracker_.sweep(); auto callback = (this->impl->enable_profiling || MIOPEN_GPU_SYNC) ? this->impl->elapsed_time_handler() : nullptr; @@ -1030,9 +935,6 @@ Handle::CreateSubBuffer(ConstData_t data, std::size_t offset, std::size_t size) const rocblas_handle_ptr& Handle::rhandle() const { - if(miopenHandle_exclusive_stream != nullptr) - MIOPEN_THROW("rocBLAS handles are bound to the stream pool and cannot be used while an " - "exclusive stream is set"); if(miopenHandle_current_stream_id == 0) return this->impl->rhandle_; // locking only if handle in multistream mode @@ -1053,9 +955,6 @@ rocblas_handle_ptr Handle::CreateRocblasHandle(miopenAcceleratorQueue_t stream) #if MIOPEN_USE_HIPBLASLT const hipblasLt_handle_ptr& Handle::HipblasLtHandle() const { - if(miopenHandle_exclusive_stream != nullptr) - MIOPEN_THROW("hipBLASLt handles are bound to the stream pool and cannot be used while an " - "exclusive stream is set"); if(miopenHandle_current_stream_id == 0) return this->impl->hip_blasLt_handle; // locking only if handle in multistream mode diff --git a/projects/miopen/src/include/miopen/any_solver.hpp b/projects/miopen/src/include/miopen/any_solver.hpp index f20b40bc6fff..3b93d0837bf4 100644 --- a/projects/miopen/src/include/miopen/any_solver.hpp +++ b/projects/miopen/src/include/miopen/any_solver.hpp @@ -90,6 +90,15 @@ struct AnySolver assert(ptr_value != nullptr); return ptr_value->GetWti(ctx, problem); }; + /// How far off the pace this solver expects to be here -- see SolverSpeedClass. + /// Exposed through the type-erased interface so the immediate-mode fallback, which + /// only holds solver::Id, can apply the same last-resort policy Find applies. + SolverSpeedClass GetSpeedClass(const ExecutionContext& ctx, + const miopen::conv::ProblemDescription& problem) const + { + assert(ptr_value != nullptr); + return ptr_value->GetSpeedClass(ctx, problem); + }; const std::type_info& Type() const { assert(ptr_value != nullptr); @@ -192,23 +201,26 @@ struct AnySolver virtual bool IsDynamic() const = 0; virtual float GetWti(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem) const = 0; - virtual const std::type_info& Type() const = 0; - virtual std::string GetSolverDbId() const = 0; + virtual SolverSpeedClass + GetSpeedClass(const ExecutionContext& ctx, + const miopen::conv::ProblemDescription& problem) const = 0; + virtual const std::type_info& Type() const = 0; + virtual std::string GetSolverDbId() const = 0; virtual ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, PerformanceDb& db, const miopen::AnyInvokeParams& invoke_ctx, - const std::string& perf_cfg) const = 0; + const std::string& perf_cfg) const = 0; virtual ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, std::function& db_getter, const miopen::AnyInvokeParams& invoke_ctx, - const std::string& perf_cfg) const = 0; + const std::string& perf_cfg) const = 0; virtual ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, DbGetter& db_getter, const miopen::AnyInvokeParams& invoke_ctx, - const std::string& perf_cfg) const = 0; + const std::string& perf_cfg) const = 0; virtual std::string GenericSearch(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, @@ -413,6 +425,12 @@ struct AnySolver { return value.GetWti(ctx, problem); } + SolverSpeedClass + GetSpeedClass(const ExecutionContext& ctx, + const miopen::conv::ProblemDescription& problem) const override + { + return value.GetSpeedClass(ctx, problem); + } ConvSolution FindSolution(const ExecutionContext& ctx, const miopen::conv::ProblemDescription& problem, diff --git a/projects/miopen/src/include/miopen/conv/solver_finders.hpp b/projects/miopen/src/include/miopen/conv/solver_finders.hpp index ff0466977e9f..4b32fb3ebfab 100644 --- a/projects/miopen/src/include/miopen/conv/solver_finders.hpp +++ b/projects/miopen/src/include/miopen/conv/solver_finders.hpp @@ -32,9 +32,13 @@ #include #include #include +#include #include +#include #include +#include +#include #include #include #include @@ -164,16 +168,30 @@ struct FindCoreResult bool is_optimal; }; -std::vector EvaluateInvokers(const Handle& handle, - const std::vector& solutions, - const AlgorithmName& algorithm_name, - const NetworkConfig& network_config, - const AnyInvokeParams& invoke_ctx, - FindCoreResult& core_result, - bool force_attach_binary, - bool& non_naive_succeeded, - bool non_naive_exists = false, - bool naive_exceeds_work = false); +/// Whether the solver speed-class policy is active (MIOPEN_DEBUG_DEFER_SLOW_SOLVERS, +/// default on). Exposed so Find and immediate mode consult one definition; when it is off +/// both behave as though SolverSpeedClass did not exist. +MIOPEN_INTERNALS_EXPORT bool IsDeferSlowSolversEnabled(); + +std::vector +EvaluateInvokers(const Handle& handle, + const std::vector& solutions, + const AlgorithmName& algorithm_name, + const NetworkConfig& network_config, + const AnyInvokeParams& invoke_ctx, + FindCoreResult& core_result, + bool force_attach_binary, + /// What each solver reported from + /// SolverInterface::GetSpeedClass for this problem. Only + /// non-Normal entries need to be present; anything absent + /// is treated as Normal, so an empty map disables the + /// policy entirely. + const std::map& speed_classes = {}, + /// The worst class to benchmark: solvers in a worse class + /// are skipped. Callers pass the best class that has an + /// applicable member, so a sole candidate is never + /// dropped however slow it declares itself to be. + solver::SolverSpeedClass benchmark_class = solver::SolverSpeedClass::Normal); FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, const ExecutionContext& ctx, diff --git a/projects/miopen/src/include/miopen/conv/solvers.hpp b/projects/miopen/src/include/miopen/conv/solvers.hpp index b690b32c9409..4b74107e7782 100644 --- a/projects/miopen/src/include/miopen/conv/solvers.hpp +++ b/projects/miopen/src/include/miopen/conv/solvers.hpp @@ -2307,6 +2307,15 @@ struct MIOPEN_INTERNALS_EXPORT ConvDirectNaiveConvFwd final : ConvSolver { return 0.01f; } + /// Naive stays *applicable* at any problem size -- it is MIOpen's universal + /// fallback and must never stop being available. But its kernel is un-tiled, so + /// past a certain amount of work a single launch runs for seconds and trips the + /// OS GPU watchdog (a TDR) merely by being benchmarked. Reporting + /// SolverSpeedClass::ExceedsLaunchBudget here lets Find and immediate mode demote it + /// to a last resort -- behind even a merely slow solver -- without touching + /// applicability. @see ConvDirectNaiveConvExceedsWorkLimit + SolverSpeedClass GetSpeedClass(const ExecutionContext&, + const miopen::conv::ProblemDescription&) const override; ConvSolution GetSolution(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; }; @@ -2327,6 +2336,9 @@ struct MIOPEN_INTERNALS_EXPORT ConvDirectNaiveConvBwd final : ConvSolver { return 0.01f; } + /// @see ConvDirectNaiveConvFwd::GetSpeedClass + SolverSpeedClass GetSpeedClass(const ExecutionContext&, + const miopen::conv::ProblemDescription&) const override; ConvSolution GetSolution(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; }; @@ -2347,6 +2359,9 @@ struct MIOPEN_INTERNALS_EXPORT ConvDirectNaiveConvWrw final : ConvSolver { return 0.01f; } + /// @see ConvDirectNaiveConvFwd::GetSpeedClass + SolverSpeedClass GetSpeedClass(const ExecutionContext&, + const miopen::conv::ProblemDescription&) const override; ConvSolution GetSolution(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; }; @@ -2375,8 +2390,8 @@ struct MIOPEN_INTERNALS_EXPORT GemmFwd1x1_0_2 final : GemmFwdBase bool MayNeedWorkspace() const override { return true; } - bool IsSlow(const ExecutionContext& context, - const miopen::conv::ProblemDescription& problem) const override; + SolverSpeedClass GetSpeedClass(const ExecutionContext& context, + const miopen::conv::ProblemDescription& problem) const override; bool IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; @@ -2414,8 +2429,8 @@ struct MIOPEN_INTERNALS_EXPORT GemmFwd1x1_0_1 final : GemmFwdBase bool MayNeedWorkspace() const override { return true; } - bool IsSlow(const ExecutionContext& context, - const miopen::conv::ProblemDescription& problem) const override; + SolverSpeedClass GetSpeedClass(const ExecutionContext& context, + const miopen::conv::ProblemDescription& problem) const override; bool IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; @@ -2435,8 +2450,8 @@ struct MIOPEN_INTERNALS_EXPORT GemmFwdRest final : GemmFwdBase bool MayNeedWorkspace() const override { return true; } - bool IsSlow(const ExecutionContext& context, - const miopen::conv::ProblemDescription& problem) const override; + SolverSpeedClass GetSpeedClass(const ExecutionContext& context, + const miopen::conv::ProblemDescription& problem) const override; bool IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; @@ -2468,8 +2483,8 @@ struct MIOPEN_INTERNALS_EXPORT GemmBwd1x1_stride2 final : GemmBwdBase bool MayNeedWorkspace() const override { return true; } - bool IsSlow(const ExecutionContext& context, - const miopen::conv::ProblemDescription& problem) const override; + SolverSpeedClass GetSpeedClass(const ExecutionContext& context, + const miopen::conv::ProblemDescription& problem) const override; bool IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; @@ -2489,8 +2504,8 @@ struct MIOPEN_INTERNALS_EXPORT GemmBwd1x1_stride1 final : GemmBwdBase bool MayNeedWorkspace() const override { return true; } - bool IsSlow(const ExecutionContext& context, - const miopen::conv::ProblemDescription& problem) const override; + SolverSpeedClass GetSpeedClass(const ExecutionContext& context, + const miopen::conv::ProblemDescription& problem) const override; bool IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription& problem) const override; @@ -2510,8 +2525,8 @@ struct MIOPEN_INTERNALS_EXPORT GemmBwdRest final : GemmBwdBase bool MayNeedWorkspace() const override { return true; } - bool IsSlow(const ExecutionContext& context, - const miopen::conv::ProblemDescription& problem) const override; + SolverSpeedClass GetSpeedClass(const ExecutionContext& context, + const miopen::conv::ProblemDescription& problem) const override; bool IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; @@ -2537,8 +2552,8 @@ struct MIOPEN_INTERNALS_EXPORT GemmWrw1x1_stride1 final : GemmWrwBase { const std::string& SolverDbId() const override { return GetSolverDbId(); } - bool IsSlow(const ExecutionContext& context, - const miopen::conv::ProblemDescription& problem) const override; + SolverSpeedClass GetSpeedClass(const ExecutionContext& context, + const miopen::conv::ProblemDescription& problem) const override; bool IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; @@ -2563,8 +2578,8 @@ struct MIOPEN_INTERNALS_EXPORT GemmWrwUniversal final : GemmWrwBase bool MayNeedWorkspace() const override { return true; } - bool IsSlow(const ExecutionContext& context, - const miopen::conv::ProblemDescription& problem) const override; + SolverSpeedClass GetSpeedClass(const ExecutionContext& context, + const miopen::conv::ProblemDescription& problem) const override; bool IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; diff --git a/projects/miopen/src/include/miopen/find_solution.hpp b/projects/miopen/src/include/miopen/find_solution.hpp index b8458dd594ff..f102a97e318d 100644 --- a/projects/miopen/src/include/miopen/find_solution.hpp +++ b/projects/miopen/src/include/miopen/find_solution.hpp @@ -416,10 +416,13 @@ struct SolverContainer { MIOPEN_LOG_I2(solver.SolverDbId() << ": Not applicable"); } - else if(env::enabled(MIOPEN_SEARCH_CUTOFF) && solver.IsSlow(ctx, problem)) - { - MIOPEN_LOG_I2(solver.SolverDbId() << ": Skipped (slow, search cutoff active)"); - } + // No GetSpeedClass() check here on purpose. A solver reporting a non-Normal + // speed class is still applicable and must stay in the candidate list, + // because it may be the *only* candidate -- dropping it here would turn a + // slow convolution into a failed one. The last-resort policy needs the + // cross-solver view that only the consumers have, so it lives in + // FindCore/EvaluateInvokers (Find) and GetSolutionsFallback (immediate + // mode). @see ShouldSkipSlowBenchmark else { const Solution s = @@ -475,10 +478,8 @@ struct SolverContainer { MIOPEN_LOG_I2(solver.SolverDbId() << ": Not applicable"); } - else if(env::enabled(MIOPEN_SEARCH_CUTOFF) && solver.IsSlow(ctx, problem)) - { - MIOPEN_LOG_I2(solver.SolverDbId() << ": Skipped (slow, search cutoff active)"); - } + // No GetSpeedClass() check here on purpose -- see the note in the sibling + // loop above. else { auto db = [&]() -> PerformanceDb& { diff --git a/projects/miopen/src/include/miopen/handle.hpp b/projects/miopen/src/include/miopen/handle.hpp index 9b6694bf2fe6..6cbd3f7ffd5c 100644 --- a/projects/miopen/src/include/miopen/handle.hpp +++ b/projects/miopen/src/include/miopen/handle.hpp @@ -38,7 +38,6 @@ #include #include #include -#include #include #include @@ -95,19 +94,6 @@ struct MIOPEN_EXPORT Handle : miopenHandle void SetStreamFromPool(int streamID) const; void ReserveExtraStreamsInPool(int cnt) const; - /// Creates a stream that belongs to the caller rather than to this Handle's - /// index-addressed stream pool. - StreamTracker::StreamPtr CreateExclusiveStream() const; - - /// Redirects GetStream() on the calling thread to \p stream, overriding the - /// stream-pool index. Pass nullptr to restore the index-selected stream. - /// Only valid for work that never reaches rocBLAS or hipBLASLt, whose - /// handles remain bound to the pool stream and will throw while an - /// exclusive stream is set. - void SetExclusiveStream(miopenAcceleratorQueue_t stream) const; - - StreamTracker& GetStreamTracker() const; - void SetAllocator(miopenAllocatorFunction allocator, miopenDeallocatorFunction deallocator, void* allocatorContext) const; @@ -198,7 +184,6 @@ struct MIOPEN_EXPORT Handle : miopenHandle void Copy(ConstData_t src, Data_t dest, std::size_t size) const; Allocator::ManageDataPtr Create(std::size_t sz) const; - std::shared_ptr GetScratchBuffer(std::size_t sz) const; Allocator::ManageDataPtr& WriteTo(const void* data, Allocator::ManageDataPtr& ddata, std::size_t sz) const; void WriteTo(const void* data, Data_t ddata, std::size_t sz) const; diff --git a/projects/miopen/src/include/miopen/solver.hpp b/projects/miopen/src/include/miopen/solver.hpp index 3d72fcc78317..67c2699eb683 100644 --- a/projects/miopen/src/include/miopen/solver.hpp +++ b/projects/miopen/src/include/miopen/solver.hpp @@ -122,6 +122,42 @@ struct SolverBase } }; +/// How far off the pace a solver expects to be for a given problem on a given device -- +/// and, at the worst level, whether a single launch is safe to *benchmark* at all. +/// +/// Some solvers are deliberately applicable far outside the region where they are a +/// sensible choice, because something has to stay applicable everywhere. Find times every +/// applicable solver by executing it, so those solvers cost real time merely by being +/// candidates. This enum lets a solver say so, and lets the framework rank the answers. +/// +/// The levels are ordered: lower is better, and the framework benchmarks the best level +/// that has any applicable member and skips the rest. Consequences worth stating, since +/// the whole policy follows from that one rule: +/// * Any Normal candidate present -> behaviour is exactly as if this did not exist. +/// * No Normal, several Slow -> all the Slow ones are benchmarked and the least slow +/// wins on measurement, which is the best available answer. +/// * Only ExceedsLaunchBudget left -> it is benchmarked, because the alternative is +/// failing a convolution that MIOpen promises to serve. A huge sole-naive shape may +/// then still trip the watchdog: an honest "extend coverage here" signal we +/// deliberately do not mask. +/// The level that must never be crossed implicitly is Slow -> ExceedsLaunchBudget: a +/// merely wasteful solver is always preferable to one that may reset the display driver. +enum class SolverSpeedClass +{ + /// A normal candidate. Nothing to defer for. + Normal = 0, + /// Applicable, and safe to launch, but expected to be far off the pace for this + /// problem -- e.g. explicit GEMM in shapes where im2col expansion and per-batch BLAS + /// calls dominate. Benchmarking one wastes Find time; it does not endanger the device. + Slow = 1, + /// A *single* launch is estimated to outlast the OS GPU watchdog, i.e. benchmarking + /// this solver risks a TDR / driver reset -- e.g. the un-tiled naive kernel on a large + /// problem. Bounding how long Find *waits* cannot avert that: the dispatch has already + /// happened and the kernel keeps occupying the GPU after the wait is abandoned, which + /// is exactly what the watchdog measures. Only a pre-launch gate works. + ExceedsLaunchBudget = 2, +}; + #ifdef _WIN32 // Suppress -Wundefined-func-template warning for Windows #pragma GCC diagnostic push @@ -152,8 +188,12 @@ struct SolverInterface : SolverBase /// Returns the workspace size required by the solver for the given Problem virtual size_t GetWorkspaceSize(const Context&, const Problem&) const { return 0; }; - /// Returns true if the solver is expected to be slow for the given problem. - virtual bool IsSlow(const Context&, const Problem&) const { return false; }; + /// Returns how far off the pace this solver expects to be for the given problem, and + /// whether a single launch is safe to benchmark. @see SolverSpeedClass. + virtual SolverSpeedClass GetSpeedClass(const Context&, const Problem&) const + { + return SolverSpeedClass::Normal; + }; }; /// Common interface for non-tunable solvers diff --git a/projects/miopen/src/include/miopen/solver/conv_direct_naive_conv.hpp b/projects/miopen/src/include/miopen/solver/conv_direct_naive_conv.hpp index 3a15523dd450..c42e729eb39a 100644 --- a/projects/miopen/src/include/miopen/solver/conv_direct_naive_conv.hpp +++ b/projects/miopen/src/include/miopen/solver/conv_direct_naive_conv.hpp @@ -52,16 +52,31 @@ std::string ConvDirectNaiveConvCompileOption(const ExecutionContext& ctx, bool ConvDirectNaiveConvIsApplicableByKernelType(const ExecutionContext&, const miopen::conv::ProblemDescription&); -/// Returns true when the problem's total MAC work exceeds the naive-conv work -/// limit. Naive stays applicable at any size (it is the universal fallback); this -/// is consumed at the selection layer (EvaluateInvokers) to keep the un-tiled -/// naive kernel from being benchmarked when it would run long enough to trip the -/// OS GPU watchdog / TDR and a non-naive alternative also applies. -/// Overridable via MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK; see the .cpp for detail. -/// Exported (MIOPEN_INTERNALS_EXPORT) so the gtest unit test can link against it -/// directly and exercise the metric on the CPU without launching a kernel. +/// Total multiply-accumulate count for the problem: the naive kernel is un-tiled, +/// so its wall-time is proportional to this. +MIOPEN_INTERNALS_EXPORT size_t ConvDirectNaiveConvWork(const miopen::conv::ProblemDescription&); + +/// Work budget (in MACs) a single naive launch may spend on a device with the given +/// capabilities, derived from a wall-time budget so the limit scales with the part +/// actually running the kernel. Takes plain integers rather than a Handle so the +/// model can be unit tested on the CPU against synthetic devices. +/// Tunable via MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS (budget) and +/// MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK (absolute override); see the .cpp. +MIOPEN_INTERNALS_EXPORT size_t ConvDirectNaiveConvWorkLimit(size_t hw_cus, + size_t wavefront_width, + size_t clock_khz); + +/// Returns true when the problem's total MAC work exceeds what the context's device +/// can retire inside the wall-time budget. Naive stays applicable at any size (it is +/// the universal fallback); this is consumed at the selection layers +/// (EvaluateInvokers for Find, GetSolutionsFallback for immediate mode) to drop the +/// un-tiled naive kernel when it would run long enough to trip the OS GPU watchdog / +/// TDR and a non-naive alternative also applies. +/// Exported (MIOPEN_INTERNALS_EXPORT) so the gtest unit tests can link against these +/// directly and exercise the model on the CPU without launching a kernel. MIOPEN_INTERNALS_EXPORT bool -ConvDirectNaiveConvExceedsWorkLimit(const miopen::conv::ProblemDescription&); +ConvDirectNaiveConvExceedsWorkLimit(const ExecutionContext&, + const miopen::conv::ProblemDescription&); bool IsInputFp32(const miopen::conv::ProblemDescription&); bool IsInputFp16(const miopen::conv::ProblemDescription&); diff --git a/projects/miopen/src/include/miopen/stream_tracker.hpp b/projects/miopen/src/include/miopen/stream_tracker.hpp deleted file mode 100644 index 3d36897a589e..000000000000 --- a/projects/miopen/src/include/miopen/stream_tracker.hpp +++ /dev/null @@ -1,74 +0,0 @@ -// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT - -#ifndef GUARD_MIOPEN_STREAM_TRACKER_HPP_ -#define GUARD_MIOPEN_STREAM_TRACKER_HPP_ - -#include -#include - -#include -#include -#include -#include - -namespace miopen { - -struct Handle; - -struct ScratchAllocation -{ - Allocator::ManageDataPtr buffer; - std::size_t size = 0; -}; - -/// Pool of streams dedicated to speculative kernel evaluation. The streams are -/// created and owned here rather than taken from the Handle's index-addressed -/// stream pool, whose ids MHA and RNN claim by hardcoded number: a slot handed -/// out here must not be one another solver is also writing to. -struct MIOPEN_INTERNALS_EXPORT StreamTracker -{ - using StreamPtr = std::shared_ptr>; - - struct Slot - { - hipStream_t stream = nullptr; - std::shared_ptr scratch; - }; - - StreamTracker() = default; - - /// Blocks until every abandoned stream has drained. Kernels left running by - /// a timed-out evaluation execute from code objects the Handle unloads during - /// teardown, so they must finish before this tracker's owner goes away. - ~StreamTracker(); - - StreamTracker(const StreamTracker&) = delete; - StreamTracker& operator=(const StreamTracker&) = delete; - - Slot acquire(const Handle& handle); - - /// Reclaims abandoned slots whose stream has gone idle, dropping the scratch - /// references they hold. Non-blocking: a slot whose stream is still busy is - /// left in place for a later sweep. - void sweep(); - - void release(Slot slot) - { - slot.scratch.reset(); - available_.push_back(std::move(slot)); - } - - void abandon(Slot slot) { draining_.push_back(std::move(slot)); } - -private: - /// Declared first so it is destroyed last: the slots below borrow these - /// streams, and the destructor synchronizes on them before they go away. - std::vector owned_streams_; - std::vector available_; - std::vector draining_; -}; - -} // namespace miopen - -#endif // GUARD_MIOPEN_STREAM_TRACKER_HPP_ diff --git a/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp b/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp index 89a2c779a296..82bed209bd62 100644 --- a/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp +++ b/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp @@ -37,11 +37,19 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_USE_PACKED_KERNELS); -// Runtime override (in MACs) for the naive-conv work-size applicability gate. -// 0 (unset) => use the built-in NAIVE_CONV_MAX_WORK default; set to a very large -// value to effectively disable the gate. See ConvDirectNaiveConvExceedsWorkLimit. +// Absolute runtime override (in MACs) for the naive-conv work-size gate. 0 (unset) +// => derive the limit from the device's capabilities; set to a very large value to +// effectively disable the gate. Wins over the time budget below. +// See ConvDirectNaiveConvWorkLimit. MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK); +// Wall-time budget (in milliseconds) that one naive launch may occupy. This is the +// quantity the gate actually cares about -- it is converted to a MAC budget using +// the live device's throughput. 0 (unset) => NAIVE_CONV_DEFAULT_BUDGET_MS. Lower it +// on a short-watchdog OS, raise it when the watchdog is disabled or TdrDelay was +// extended. Ignored when MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK is set. +MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS); + namespace miopen { namespace solver { @@ -315,25 +323,19 @@ bool ConvDirectNaiveConvIsApplicableByKernelType(const ExecutionContext& ctx, return true; } -// Upper bound (in multiply-accumulate operations) on the total work above which the -// naive convolution solver should not be *benchmarked* during Find. The naive -// kernel is un-tiled and the GPU saturates, so its wall-time is proportional to the -// total MAC count: N * K * C_per_group * output_spatial_volume * filter_volume. On -// very large problems (e.g. VAE-decode / 3D shapes) a single naive launch runs for -// multiple seconds; when MIOpen's Find step *benchmarks* candidate solvers it will -// execute naive and trip the OS GPU watchdog (TDR) -- even when a fast solver also -// applies and would ultimately win. Naive stays *applicable* at any size so it can -// serve as the universal fallback; the selection layer (EvaluateInvokers) consumes -// this threshold to skip benchmarking naive when it is over the limit AND a -// non-naive alternative also applies. When naive is the sole applicable solver it -// still runs, so a shape only naive can serve is served (and, if huge, may TDR on a -// short-watchdog OS -- an honest "extend coverage here" signal, not masked). -// Overridable via MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK. -// NOTE: the 16 GMAC default is a conservative estimate, not a measured value; it -// should be calibrated once on stable hardware (see project_naive_find_tdr). -constexpr size_t NAIVE_CONV_MAX_WORK = static_cast(16) * 1000 * 1000 * 1000; // ~16 GMAC - -bool ConvDirectNaiveConvExceedsWorkLimit(const ProblemDescription& problem) +// The naive kernel is un-tiled and saturates the GPU, so its wall-time is +// proportional to the total MAC count: N * K * C_per_group * output_spatial_volume * +// filter_volume. On very large problems (e.g. VAE-decode / 3D shapes) a single naive +// launch runs for multiple seconds; when MIOpen *benchmarks* candidate solvers during +// Find, or immediate mode ranks one without measuring it, executing naive trips the +// OS GPU watchdog (TDR) -- even when a fast solver also applies and would ultimately +// win. Naive stays *applicable* at any size so it can serve as the universal +// fallback; the selection layers (EvaluateInvokers and GetSolutionsFallback) consume +// this threshold to drop naive when it is over the limit AND a non-naive alternative +// also applies. When naive is the sole applicable solver it still runs, so a shape +// only naive can serve is served (and, if huge, may TDR on a short-watchdog OS -- an +// honest "extend coverage here" signal, not masked). +size_t ConvDirectNaiveConvWork(const ProblemDescription& problem) { const size_t n = problem.GetInBatchSize(); const size_t k = problem.GetOutChannels(); @@ -346,17 +348,119 @@ bool ConvDirectNaiveConvExceedsWorkLimit(const ProblemDescription& problem) const size_t fil_vol = problem.GetWeightsWidth() * problem.GetWeightsHeight() * (problem.Is2d() ? size_t{1} : problem.GetWeightsDepth()); - const size_t work = n * k * c_per_group * out_vol * fil_vol; + return n * k * c_per_group * out_vol * fil_vol; +} + +// Sustained MACs the *naive* kernel retires per lane per 1000 clocks. The naive inner +// loop is a single FMA fed by two un-cached global loads, with no LDS staging and no +// register blocking, so it is memory-latency bound and runs far below the +// one-FMA-per-lane-per-clock issue peak (which in these units would be 1000). +// +// Calibrated against two measured sweeps totalling 109 timed shape/direction points: +// gfx1151 (RDNA3.5, 80 hw CUs, wave32, 2.9 GHz) and gfx950 (CDNA4, 256 hw CUs, wave64, +// 2.4 GHz). Over the 39 shape/direction pairs measured on both, gfx950 ran a median +// 5.04x faster on shapes >= 5 GMAC against a 5.30x lane-cycle ratio -- so the linear +// device scaling this model assumes holds to within ~5% across a 5x spread of part size +// and across both wavefront widths. That scaling, not the constant, is the load-bearing +// claim; it is what lets one number cover CDNA and RDNA. +// +// 3 is the conservative end of the measured spread rather than its centre: implied rates +// centre on a median of 11 on gfx1151 and 7 on gfx950, so the derived limit fires well +// before naive reaches the budget on a typical shape. It also keeps the derived limits +// in the same magnitude as the fixed 16 GMAC this model replaces, so the change +// contributes device *scaling* without simultaneously moving the absolute calibration. +// Too high risks a TDR; too low over-gates naive on shapes it would have won (the class +// of regression #9513 was filed against). +// +// Known weakness, deliberately not addressed here: MAC count is not occupancy-aware. +// Naive's grid is group * n * k_per_group blocks wide, so a shape whose (n, k_per_group) +// product is below the part's CU count under-fills it and runs far under the modeled +// throughput. Every measured point that fell short of the model is of this form -- 3D +// convolutions at n=1, k=64 on a 256-CU part -- and one of them (32x128x128, bwd, +// 58 GMAC) reached 1256 ms against the 1000 ms budget. Such shapes are *under*-gated: +// safe for coverage, unsafe for TDR. Note that measurement predates the intra-grid +// spatial tiling added in #7529, which widens exactly these grids. +// +// To re-measure: force naive with MIOPEN_DEBUG_FIND_ONLY_SOLVER, pin +// MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK high so the gate does not skip what you are +// timing, and run MIOpenDriver -t 1 over a range of shapes. Each shape then implies +// work_macs * 1000 / (hw_cus * wavefront * clock_khz * measured_ms), which is this +// constant. Note hw_cus is the doubled value from Handle::GetMaxHardwareComputeUnits() +// on gfx1*. The spread of that quantity across shapes also indicates whether MAC count +// is predictive at all, which is the assumption the whole gate rests on. +constexpr size_t NAIVE_MAC_PER_1K_LANE_CYCLES = 3; + +// Wall-time a single naive launch may occupy, default and ceiling. Windows' default +// TdrDelay is 2 s, so 1 s leaves 2x margin. The ceiling exists only to keep an +// outlandish env value from overflowing the multiply below. +constexpr size_t NAIVE_CONV_DEFAULT_BUDGET_MS = 1000; +constexpr size_t NAIVE_CONV_MAX_BUDGET_MS = 60 * 1000; + +// Limit used when the device cannot report the capabilities the model needs. This is +// the fixed pre-device-aware default; it is conservative for every current part. +constexpr size_t NAIVE_CONV_FALLBACK_MAX_WORK = + static_cast(16) * 1000 * 1000 * 1000; // ~16 GMAC + +// Convert a wall-time budget into a MAC budget using the live device's throughput. +// What we need to bound is the *duration* of one naive launch -- on a short-watchdog +// OS, a launch that outlives the watchdog is a GPU reset, and by the time it fires it +// is far too late to recover. We cannot measure the launch without running it, so we +// bound the work instead: +// +// lane_cycles/ms = hw_cus * wavefront_width * clock_khz +// naive_macs/ms = lane_cycles/ms * NAIVE_MAC_PER_1K_LANE_CYCLES / 1000 +// limit (MACs) = naive_macs/ms * budget_ms +// +// Clock arrives in kHz, which is already "cycles per ms", so the whole computation +// stays in integer math with no unit juggling. Every term but the throughput constant +// is read from the device, so the limit tracks the part actually running the kernel -- +// a 304-CU MI300X earns ~123 GMAC at the default budget while a 32-WGP iGPU earns +// ~8.9 GMAC, with no per-arch table to hand-maintain. +// +// Takes capabilities as plain integers rather than a Handle so the model is unit +// testable on the CPU against synthetic devices. +size_t ConvDirectNaiveConvWorkLimit(size_t hw_cus, size_t wavefront_width, size_t clock_khz) +{ + const size_t absolute_override = env::value(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK); + if(absolute_override != 0) + return absolute_override; + + // A device that cannot report its own capabilities gives the model nothing to + // scale from. Fall back to the fixed limit rather than deriving a nonsense + // (possibly zero) budget that would gate every shape. + if(hw_cus == 0 || wavefront_width == 0 || clock_khz == 0) + { + MIOPEN_LOG_I2("ConvDirectNaiveConv work limit: incomplete device capabilities (cus " + << hw_cus << ", wavefront " << wavefront_width << ", clock " << clock_khz + << " kHz); using fixed fallback " << NAIVE_CONV_FALLBACK_MAX_WORK << " MACs"); + return NAIVE_CONV_FALLBACK_MAX_WORK; + } - size_t limit = env::value(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK); - if(limit == 0) - limit = NAIVE_CONV_MAX_WORK; + size_t budget_ms = env::value(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS); + if(budget_ms == 0) + budget_ms = NAIVE_CONV_DEFAULT_BUDGET_MS; + budget_ms = std::min(budget_ms, NAIVE_CONV_MAX_BUDGET_MS); + + // clock_khz is cycles/ms, so this product is lane-cycles/ms directly. + const size_t macs_per_ms = + hw_cus * wavefront_width * clock_khz * NAIVE_MAC_PER_1K_LANE_CYCLES / 1000; + return macs_per_ms * budget_ms; +} + +bool ConvDirectNaiveConvExceedsWorkLimit(const ExecutionContext& ctx, + const ProblemDescription& problem) +{ + const auto& handle = ctx.GetStream(); + const size_t work = ConvDirectNaiveConvWork(problem); + const size_t limit = ConvDirectNaiveConvWorkLimit( + handle.GetMaxHardwareComputeUnits(), handle.GetWavefrontWidth(), handle.GetClockRateKhz()); if(work > limit) { MIOPEN_LOG_I2("ConvDirectNaiveConv over work limit: work " - << work << " MACs exceeds limit " << limit - << " -- defer from Find benchmark when an alternative applies (avoids " + << work << " MACs exceeds limit " << limit << " for " + << handle.GetDeviceName() + << " -- defer from selection when an alternative applies (avoids " "TDR-prone naive launch)"); return true; } diff --git a/projects/miopen/src/solver/conv/conv_direct_naive_conv_bwd.cpp b/projects/miopen/src/solver/conv/conv_direct_naive_conv_bwd.cpp index 550f140cc917..6db993365b8e 100644 --- a/projects/miopen/src/solver/conv/conv_direct_naive_conv_bwd.cpp +++ b/projects/miopen/src/solver/conv/conv_direct_naive_conv_bwd.cpp @@ -79,6 +79,13 @@ bool ConvDirectNaiveConvBwd::IsApplicable(const ExecutionContext& ctx, return true; } +SolverSpeedClass ConvDirectNaiveConvBwd::GetSpeedClass(const ExecutionContext& ctx, + const ProblemDescription& problem) const +{ + return ConvDirectNaiveConvExceedsWorkLimit(ctx, problem) ? SolverSpeedClass::ExceedsLaunchBudget + : SolverSpeedClass::Normal; +} + ConvSolution ConvDirectNaiveConvBwd::GetSolution(const ExecutionContext& ctx, const ProblemDescription& problem) const { diff --git a/projects/miopen/src/solver/conv/conv_direct_naive_conv_fwd.cpp b/projects/miopen/src/solver/conv/conv_direct_naive_conv_fwd.cpp index 67e0bb66b03a..666b6d500f12 100644 --- a/projects/miopen/src/solver/conv/conv_direct_naive_conv_fwd.cpp +++ b/projects/miopen/src/solver/conv/conv_direct_naive_conv_fwd.cpp @@ -81,6 +81,13 @@ bool ConvDirectNaiveConvFwd::IsApplicable(const ExecutionContext& ctx, return true; } +SolverSpeedClass ConvDirectNaiveConvFwd::GetSpeedClass(const ExecutionContext& ctx, + const ProblemDescription& problem) const +{ + return ConvDirectNaiveConvExceedsWorkLimit(ctx, problem) ? SolverSpeedClass::ExceedsLaunchBudget + : SolverSpeedClass::Normal; +} + ConvSolution ConvDirectNaiveConvFwd::GetSolution(const ExecutionContext& ctx, const ProblemDescription& problem) const { diff --git a/projects/miopen/src/solver/conv/conv_direct_naive_conv_wrw.cpp b/projects/miopen/src/solver/conv/conv_direct_naive_conv_wrw.cpp index 27767cbebbc5..481c4413594f 100644 --- a/projects/miopen/src/solver/conv/conv_direct_naive_conv_wrw.cpp +++ b/projects/miopen/src/solver/conv/conv_direct_naive_conv_wrw.cpp @@ -80,6 +80,13 @@ bool ConvDirectNaiveConvWrw::IsApplicable(const ExecutionContext& ctx, return true; } +SolverSpeedClass ConvDirectNaiveConvWrw::GetSpeedClass(const ExecutionContext& ctx, + const ProblemDescription& problem) const +{ + return ConvDirectNaiveConvExceedsWorkLimit(ctx, problem) ? SolverSpeedClass::ExceedsLaunchBudget + : SolverSpeedClass::Normal; +} + ConvSolution ConvDirectNaiveConvWrw::GetSolution(const ExecutionContext& ctx, const ProblemDescription& problem) const { diff --git a/projects/miopen/src/solver/conv/gemm.cpp b/projects/miopen/src/solver/conv/gemm.cpp index ea6445374959..9af040239f34 100644 --- a/projects/miopen/src/solver/conv/gemm.cpp +++ b/projects/miopen/src/solver/conv/gemm.cpp @@ -216,9 +216,11 @@ size_t GemmFwd1x1_0_2::GetWorkspaceSize(const ExecutionContext& context, #endif } -bool GemmFwd1x1_0_2::IsSlow(const ExecutionContext& context, - const ProblemDescription& problem) const +SolverSpeedClass GemmFwd1x1_0_2::GetSpeedClass(const ExecutionContext& context, + const ProblemDescription& problem) const { + // Slow, never ExceedsLaunchBudget: explicit GEMM runs in bounded per-batch chunks, so + // it is wasteful to benchmark here, not a watchdog risk. @see SolverSpeedClass const std::string& arch = context.GetStream().GetDeviceName(); const std::set mi = {"gfx942", "gfx955"}; const bool is_mi = mi.find(arch) != mi.end(); @@ -233,7 +235,7 @@ bool GemmFwd1x1_0_2::IsSlow(const ExecutionContext& context, if(is_gfx11 || is_gfx12) { - return false; + return SolverSpeedClass::Normal; } else if(is_mi) { @@ -241,10 +243,10 @@ bool GemmFwd1x1_0_2::IsSlow(const ExecutionContext& context, // SWPG < 150k: Low spatial-channel work // CPG < 48: Extremely low channels per group (critical discriminator) if(spatial_work_per_group < 150000 && channels_per_group < 48) - return true; + return SolverSpeedClass::Slow; } - return false; + return SolverSpeedClass::Normal; } bool GemmFwd1x1_0_2::IsApplicable(const ExecutionContext& context, @@ -669,9 +671,11 @@ size_t GemmFwd1x1_0_1::GetWorkspaceSize(const ExecutionContext&, const ProblemDe return 0; } -bool GemmFwd1x1_0_1::IsSlow(const ExecutionContext& context, - const ProblemDescription& problem) const +SolverSpeedClass GemmFwd1x1_0_1::GetSpeedClass(const ExecutionContext& context, + const ProblemDescription& problem) const { + // Slow, never ExceedsLaunchBudget: explicit GEMM runs in bounded per-batch chunks, so + // it is wasteful to benchmark here, not a watchdog risk. @see SolverSpeedClass const std::string& arch = context.GetStream().GetDeviceName(); const std::set mi = {"gfx942", "gfx955"}; const bool is_mi = mi.find(arch) != mi.end(); @@ -686,7 +690,7 @@ bool GemmFwd1x1_0_1::IsSlow(const ExecutionContext& context, if(is_gfx11 || is_gfx12) { - return false; + return SolverSpeedClass::Normal; } else if(is_mi) { @@ -694,10 +698,10 @@ bool GemmFwd1x1_0_1::IsSlow(const ExecutionContext& context, // SWPG < 200k: Low spatial-channel work (memory-bound) // CPG < 704: Moderate channels (poor reuse) if(spatial_work_per_group < 200000 && channels_per_group < 704) - return true; + return SolverSpeedClass::Slow; } - return false; + return SolverSpeedClass::Normal; } bool GemmFwd1x1_0_1::IsApplicable(const ExecutionContext& context, @@ -987,8 +991,11 @@ size_t GemmFwdRest::GetWorkspaceSize(const ExecutionContext& context, #endif } -bool GemmFwdRest::IsSlow(const ExecutionContext& context, const ProblemDescription& problem) const +SolverSpeedClass GemmFwdRest::GetSpeedClass(const ExecutionContext& context, + const ProblemDescription& problem) const { + // Slow, never ExceedsLaunchBudget: explicit GEMM runs in bounded per-batch chunks, so + // it is wasteful to benchmark here, not a watchdog risk. @see SolverSpeedClass const std::string& arch = context.GetStream().GetDeviceName(); const std::set mi = {"gfx942", "gfx955"}; const bool is_mi = mi.find(arch) != mi.end(); @@ -1014,7 +1021,7 @@ bool GemmFwdRest::IsSlow(const ExecutionContext& context, const ProblemDescripti // SPB < 50: Low spatial-per-batch = batch fragmentation // Performance: FPR=11-25%, TPR=61-76%, Score=1.37-1.60 if(spatial_per_batch < 50) - return true; + return SolverSpeedClass::Slow; } else if(is_mi) { @@ -1022,14 +1029,14 @@ bool GemmFwdRest::IsSlow(const ExecutionContext& context, const ProblemDescripti // SWPG < 2.5M: Low spatial-channel work (memory-bound) // CPG < 56: Very low channels (minimal reuse) if(spatial_work_per_group < 2500000 && channels_per_group < 56) - return true; + return SolverSpeedClass::Slow; // SECONDARY: Batch fragmentation detection if(spatial_per_batch < 24.0) - return true; + return SolverSpeedClass::Slow; } - return false; + return SolverSpeedClass::Normal; } bool GemmFwdRest::IsApplicable(const ExecutionContext& context, diff --git a/projects/miopen/src/solver/conv/gemm_bwd.cpp b/projects/miopen/src/solver/conv/gemm_bwd.cpp index 995b7d7b58dc..cbb7cfdc500c 100644 --- a/projects/miopen/src/solver/conv/gemm_bwd.cpp +++ b/projects/miopen/src/solver/conv/gemm_bwd.cpp @@ -194,9 +194,11 @@ size_t GemmBwd1x1_stride2::GetWorkspaceSize(const ExecutionContext& context, #endif } -bool GemmBwd1x1_stride2::IsSlow(const ExecutionContext& context, - const ProblemDescription& problem) const +SolverSpeedClass GemmBwd1x1_stride2::GetSpeedClass(const ExecutionContext& context, + const ProblemDescription& problem) const { + // Slow, never ExceedsLaunchBudget: explicit GEMM runs in bounded per-batch chunks, so + // it is wasteful to benchmark here, not a watchdog risk. @see SolverSpeedClass const std::string& arch = context.GetStream().GetDeviceName(); const std::set mi = {"gfx942", "gfx955"}; const bool is_mi = mi.find(arch) != mi.end(); @@ -211,7 +213,7 @@ bool GemmBwd1x1_stride2::IsSlow(const ExecutionContext& context, if(is_gfx11 || is_gfx12) { - return false; + return SolverSpeedClass::Normal; } else if(is_mi) { @@ -219,10 +221,10 @@ bool GemmBwd1x1_stride2::IsSlow(const ExecutionContext& context, // SWPG < 400k: Moderate spatial-channel work // CPG < 192: Low channels per group (critical discriminator) if(spatial_work_per_group < 400000 && channels_per_group < 192) - return true; + return SolverSpeedClass::Slow; } - return false; + return SolverSpeedClass::Normal; } bool GemmBwd1x1_stride2::IsApplicable(const ExecutionContext& context, @@ -430,9 +432,11 @@ size_t GemmBwd1x1_stride1::GetWorkspaceSize(const ExecutionContext&, return 0; } -bool GemmBwd1x1_stride1::IsSlow(const ExecutionContext& context, - const ProblemDescription& problem) const +SolverSpeedClass GemmBwd1x1_stride1::GetSpeedClass(const ExecutionContext& context, + const ProblemDescription& problem) const { + // Slow, never ExceedsLaunchBudget: explicit GEMM runs in bounded per-batch chunks, so + // it is wasteful to benchmark here, not a watchdog risk. @see SolverSpeedClass const std::string& arch = context.GetStream().GetDeviceName(); const std::set mi = {"gfx942", "gfx955"}; const bool is_mi = mi.find(arch) != mi.end(); @@ -447,7 +451,7 @@ bool GemmBwd1x1_stride1::IsSlow(const ExecutionContext& context, if(is_gfx11 || is_gfx12) { - return false; + return SolverSpeedClass::Normal; } else if(is_mi) { @@ -455,10 +459,10 @@ bool GemmBwd1x1_stride1::IsSlow(const ExecutionContext& context, // SWPG < 200k: Low spatial-channel work (memory-bound) // CPG < 640: Moderate channels (poor reuse) if(spatial_work_per_group < 200000 && channels_per_group < 640) - return true; + return SolverSpeedClass::Slow; } - return false; + return SolverSpeedClass::Normal; } bool GemmBwd1x1_stride1::IsApplicable(const ExecutionContext& context, @@ -697,8 +701,11 @@ size_t GemmBwdRest::GetWorkspaceSize(const ExecutionContext& context, #endif } -bool GemmBwdRest::IsSlow(const ExecutionContext& context, const ProblemDescription& problem) const +SolverSpeedClass GemmBwdRest::GetSpeedClass(const ExecutionContext& context, + const ProblemDescription& problem) const { + // Slow, never ExceedsLaunchBudget: explicit GEMM runs in bounded per-batch chunks, so + // it is wasteful to benchmark here, not a watchdog risk. @see SolverSpeedClass const std::string& arch = context.GetStream().GetDeviceName(); const std::set mi = {"gfx942", "gfx955"}; const bool is_mi = mi.find(arch) != mi.end(); @@ -722,12 +729,12 @@ bool GemmBwdRest::IsSlow(const ExecutionContext& context, const ProblemDescripti // SWPG < 1.6M: Low spatial-channel work // CPG < 360: Low channels if(spatial_work_per_group < 1600000 && channels_per_group < 360) - return true; + return SolverSpeedClass::Slow; // SECONDARY: Batch fragmentation detection // SPB < 0.8: Extreme batch fragmentation if(spatial_per_batch < 0.8) - return true; + return SolverSpeedClass::Slow; } else if(is_mi) { @@ -735,15 +742,15 @@ bool GemmBwdRest::IsSlow(const ExecutionContext& context, const ProblemDescripti // SWPG < 3M: Low spatial-channel work // CPG < 112: Very low channels if(spatial_work_per_group < 3000000 && channels_per_group < 112) - return true; + return SolverSpeedClass::Slow; // SECONDARY: Batch fragmentation detection // SPB < 40.0: Each batch item has < 40 pixels of spatial work if(spatial_per_batch < 40.0) - return true; + return SolverSpeedClass::Slow; } - return false; + return SolverSpeedClass::Normal; } bool GemmBwdRest::IsApplicable(const ExecutionContext& context, diff --git a/projects/miopen/src/solver/conv/gemm_wrw.cpp b/projects/miopen/src/solver/conv/gemm_wrw.cpp index 469060af12d0..6b2d5568057c 100644 --- a/projects/miopen/src/solver/conv/gemm_wrw.cpp +++ b/projects/miopen/src/solver/conv/gemm_wrw.cpp @@ -146,9 +146,11 @@ float GemmWrwBase::GetWti(const ExecutionContext&, const ProblemDescription& pro #endif } -bool GemmWrw1x1_stride1::IsSlow(const ExecutionContext& context, - const ProblemDescription& problem) const +SolverSpeedClass GemmWrw1x1_stride1::GetSpeedClass(const ExecutionContext& context, + const ProblemDescription& problem) const { + // Slow, never ExceedsLaunchBudget: explicit GEMM runs in bounded per-batch chunks, so + // it is wasteful to benchmark here, not a watchdog risk. @see SolverSpeedClass const std::string& arch = context.GetStream().GetDeviceName(); const std::set mi = {"gfx942", "gfx955"}; const bool is_mi = mi.find(arch) != mi.end(); @@ -177,17 +179,17 @@ bool GemmWrw1x1_stride1::IsSlow(const ExecutionContext& context, // Threshold: batch > 16 AND cpg < 1400 // Performance: FPR=3-15%, TPR=73-87%, Score=1.65-1.79 if(b > 16 && channels_per_group < 1400) - return true; + return SolverSpeedClass::Slow; } else if(is_mi) { // SPB-ONLY: Batch fragmentation detection // SPB < 48.0: Each batch item has < 48 pixels of spatial work if(spatial_per_batch < 48.0) - return true; + return SolverSpeedClass::Slow; } - return false; + return SolverSpeedClass::Normal; } #if MIOPEN_USE_GEMM @@ -490,9 +492,11 @@ size_t GemmWrwUniversal::GetWorkspaceSize(const ExecutionContext& context, #endif } -bool GemmWrwUniversal::IsSlow(const ExecutionContext& context, - const ProblemDescription& problem) const +SolverSpeedClass GemmWrwUniversal::GetSpeedClass(const ExecutionContext& context, + const ProblemDescription& problem) const { + // Slow, never ExceedsLaunchBudget: explicit GEMM runs in bounded per-batch chunks, so + // it is wasteful to benchmark here, not a watchdog risk. @see SolverSpeedClass const std::string& arch = context.GetStream().GetDeviceName(); const std::set mi = {"gfx942", "gfx955"}; const bool is_mi = mi.find(arch) != mi.end(); @@ -514,17 +518,17 @@ bool GemmWrwUniversal::IsSlow(const ExecutionContext& context, // SPB < 100: Low spatial-per-batch = batch fragmentation // Performance: FPR=19-27%, TPR=72-92%, Score=1.49-1.66 if(spatial_per_batch < 100) - return true; + return SolverSpeedClass::Slow; } else if(is_mi) { // SPB-ONLY: Batch fragmentation detection // SPB < 48.0: Each batch item has < 48 pixels of spatial work if(spatial_per_batch < 48.0) - return true; + return SolverSpeedClass::Slow; } - return false; + return SolverSpeedClass::Normal; } bool GemmWrwUniversal::IsApplicable(const ExecutionContext& context, diff --git a/projects/miopen/test/gtest/naive_timeout.cpp b/projects/miopen/test/gtest/naive_timeout.cpp deleted file mode 100644 index 32cc404ab240..000000000000 --- a/projects/miopen/test/gtest/naive_timeout.cpp +++ /dev/null @@ -1,213 +0,0 @@ -// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT - -// Tests for MIOPEN_NAIVE_TIMEOUT behavior in the find path. -// -// When MIOPEN_NAIVE_TIMEOUT=1 (default) and a non-naive solver has already -// succeeded, the naive solver's warmup run is launched on a disposable stream -// with a wall-clock budget of 3x the best non-naive time. If the warmup -// exceeds that budget the stream is abandoned and the naive solver is skipped. -// If the warmup finishes in time, naive proceeds through normal benchmarking -// and may appear in (or even win) the results. -// -// Paths through EvaluateInvokers naive-timeout logic (solver_finders.cpp): -// -// Path naive_timeout defer_naive non_naive_succeeded Outcome -// ---- ------------- ----------- ------------------- ------------------------- -// A false - - naive evaluated (flag off) -// B true false - naive evaluated (no non-naive exists) -// C true true false naive evaluated (non-naive all failed) -// D true true true naive timed — may pass or be skipped -// -// Tests cover paths A, B, and D. Path C requires fault injection. -// -// Path B uses MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvFwd to restrict -// the candidate set to the naive solver only. -// -// Tests assert presence/absence of naive solvers rather than which solver wins -// by timing, so results are not sensitive to GPU performance variation. -// -// Each GPU test uses ScopedFindDb to disable miopen::debug::testing_find_db_enabled -// so TryLoad always calls the regenerator (FindCore / EvaluateInvokers). - -#include - -#include -#include -#include -#include - -#include "get_handle.hpp" -#include "gtest_common.hpp" -#include "../tensor_holder.hpp" - -MIOPEN_LIB_ENV_VAR(MIOPEN_NAIVE_TIMEOUT) -MIOPEN_LIB_ENV_VAR(MIOPEN_NAIVE_TIMEOUT_FACTOR) - -namespace { - -struct ScopedFindDb -{ - ScopedFindDb() : prev_(miopen::debug::testing_find_db_enabled) - { - miopen::debug::testing_find_db_enabled = false; - } - ~ScopedFindDb() { miopen::debug::testing_find_db_enabled = prev_; } - - bool prev_; -}; - -std::vector RunFind(miopenHandle_t handle, size_t max_solutions = 8) -{ - tensor x{1, 64, 14, 14}; - tensor w{64, 64, 3, 3}; - - miopen::ConvolutionDescriptor conv{ - 2, miopenConvolution, miopenPaddingDefault, {1, 1}, {1, 1}, {1, 1}}; - tensor y{conv.GetForwardOutputTensor(x.desc, w.desc)}; - - miopenProblem_t problem = nullptr; - EXPECT_EQ(miopenCreateConvProblem(&problem, &conv, miopenProblemDirectionForward), - miopenStatusSuccess); - EXPECT_EQ(miopenSetProblemTensorDescriptor(problem, miopenTensorConvolutionX, &x.desc), - miopenStatusSuccess); - EXPECT_EQ(miopenSetProblemTensorDescriptor(problem, miopenTensorConvolutionW, &w.desc), - miopenStatusSuccess); - EXPECT_EQ(miopenSetProblemTensorDescriptor(problem, miopenTensorConvolutionY, &y.desc), - miopenStatusSuccess); - - std::vector solutions(max_solutions); - size_t num_found = 0; - EXPECT_EQ( - miopenFindSolutions(handle, problem, nullptr, solutions.data(), &num_found, max_solutions), - miopenStatusSuccess); - solutions.resize(num_found); - - std::vector names; - for(auto sol : solutions) - { - uint64_t solver_id = 0; - EXPECT_EQ(miopenGetSolutionSolverId(sol, &solver_id), miopenStatusSuccess); - names.push_back(miopen::solver::Id{solver_id}.ToString()); - miopenDestroySolution(sol); - } - - miopenDestroyProblem(problem); - return names; -} - -bool IsNaive(const std::string& name) { return name.find("Naive") != std::string::npos; } - -// RunFind's convolution (3x3, pad 1, stride 1, fp32, K=64) is exactly what -// ConvBinWinograd3x3U accepts, and that solver is applicable only on these -// archs. There IsWinograd3x3SupportedAndFast() sets use_winograd_only, which -// disables the Direct finder outright, so ConvDirectNaive is never a candidate -// and assertions about its presence or absence say nothing about the timeout. -bool IsWinogradOnlyArch(const miopen::Handle& handle) -{ - const auto name = handle.GetDeviceName(); - return name == "gfx900" || name == "gfx906" || name == "gfx908"; -} - -} // namespace - -// With MIOPEN_NAIVE_TIMEOUT=0 (opt-out), a naive solver must appear somewhere in the -// results, confirming that the timeout logic did not fire. -TEST(GPU_NaiveTimeout_FP32, NaivePresentWhenTimeoutDisabled) -{ - auto& handle_ref = get_handle(); - miopenHandle_t handle = &handle_ref; - - if(IsWinogradOnlyArch(handle_ref)) - GTEST_SKIP() << "Direct finder is disabled by use_winograd_only on " - << handle_ref.GetDeviceName(); - - ScopedFindDb no_cache; - ScopedEnvironment guard_naive(MIOPEN_NAIVE_TIMEOUT, false); - auto solvers = RunFind(handle); - - ASSERT_FALSE(solvers.empty()) << "miopenFindSolutions returned no results"; - const bool any_naive = std::any_of(solvers.begin(), solvers.end(), IsNaive); - std::string all_names; - for(const auto& s : solvers) - all_names += (all_names.empty() ? "" : ", ") + s; - EXPECT_TRUE(any_naive) - << "Expected at least one naive solver in results with MIOPEN_NAIVE_TIMEOUT=0" << "; got: [" - << all_names << "]"; -} - -// With MIOPEN_NAIVE_TIMEOUT=1 and MIOPEN_DEBUG_FIND_ONLY_SOLVER restricting the candidate -// set to ConvDirectNaiveConvFwd, there are no non-naive solutions so defer_naive=false and -// the timeout cannot fire (path B). Find must return a naive result rather than nothing. -TEST(GPU_NaiveTimeout_FP32, NaiveNotSuppressedWhenOnlySolver) -{ - auto& handle_ref = get_handle(); - miopenHandle_t handle = &handle_ref; - - ScopedFindDb no_cache; - ScopedEnvironment guard_naive(MIOPEN_NAIVE_TIMEOUT, true); - ScopedEnvironment only_naive(MIOPEN_DEBUG_FIND_ONLY_SOLVER, - std::string{"ConvDirectNaiveConvFwd"}); - auto solvers = RunFind(handle); - - ASSERT_FALSE(solvers.empty()) - << "miopenFindSolutions returned no results: naive must not be suppressed when it is the " - "only applicable solver (path B)"; - const bool all_naive = std::all_of(solvers.begin(), solvers.end(), IsNaive); - EXPECT_TRUE(all_naive) - << "Expected only naive solvers when MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvFwd"; -} - -// With MIOPEN_NAIVE_TIMEOUT=1 (default), find must return results. The naive solver may or -// may not appear depending on whether it finished within the 3x budget — both outcomes are -// valid. This test only verifies that the timeout mechanism does not break the find path. -TEST(GPU_NaiveTimeout_FP32, FindSucceedsWithTimeoutEnabled) -{ - auto& handle_ref = get_handle(); - miopenHandle_t handle = &handle_ref; - - ScopedFindDb no_cache; - ScopedEnvironment guard_naive(MIOPEN_NAIVE_TIMEOUT, true); - auto solvers = RunFind(handle); - - ASSERT_FALSE(solvers.empty()) << "miopenFindSolutions returned no results with timeout enabled"; -} - -// With MIOPEN_NAIVE_TIMEOUT_FACTOR=1 (0.01x budget), the wall-clock deadline expires -// before the naive kernel can finish. Naive must be absent from the results. -TEST(GPU_NaiveTimeout_FP32, NaiveSkippedWithTinyBudget) -{ - auto& handle_ref = get_handle(); - miopenHandle_t handle = &handle_ref; - - if(IsWinogradOnlyArch(handle_ref)) - GTEST_SKIP() << "Direct finder is disabled by use_winograd_only on " - << handle_ref.GetDeviceName(); - - ScopedFindDb no_cache; - ScopedEnvironment guard_naive(MIOPEN_NAIVE_TIMEOUT, true); - ScopedEnvironment guard_factor(MIOPEN_NAIVE_TIMEOUT_FACTOR, 1); - auto solvers = RunFind(handle); - - ASSERT_FALSE(solvers.empty()) << "miopenFindSolutions returned no results"; - const bool any_naive = std::any_of(solvers.begin(), solvers.end(), IsNaive); - std::string all_names; - for(const auto& s : solvers) - all_names += (all_names.empty() ? "" : ", ") + s; - EXPECT_FALSE(any_naive) << "Naive solver should have been short-circuited with 1% budget" - << "; got: [" << all_names << "]"; -} - -// CPU-only: verify that MIOPEN_NAIVE_TIMEOUT can be set and read back via the debug -// registry (the same path ScopedEnvironment uses in the GPU tests). -TEST(CPU_NaiveTimeout_NONE, EnvVarRoundTrips) -{ - { - ScopedEnvironment guard(MIOPEN_NAIVE_TIMEOUT, true); - EXPECT_TRUE(lib_env::value(MIOPEN_NAIVE_TIMEOUT)); - } - { - ScopedEnvironment guard(MIOPEN_NAIVE_TIMEOUT, false); - EXPECT_FALSE(lib_env::value(MIOPEN_NAIVE_TIMEOUT)); - } -} diff --git a/projects/miopen/test/gtest/stream_tracker.cpp b/projects/miopen/test/gtest/stream_tracker.cpp deleted file mode 100644 index ba0145c71ba6..000000000000 --- a/projects/miopen/test/gtest/stream_tracker.cpp +++ /dev/null @@ -1,304 +0,0 @@ -// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT - -#include -#include "get_handle.hpp" - -#include -#include - -#include - -#include -#include -#include - -namespace { - -struct StreamGate -{ - std::mutex mtx; - std::condition_variable cv; - bool released = false; - - static void callback(void* arg) - { - auto* self = static_cast(arg); - std::unique_lock lk(self->mtx); - self->cv.wait(lk, [self] { return self->released; }); - } - - void open() - { - { - std::lock_guard lk(mtx); - released = true; - } - cv.notify_one(); - } -}; - -class GPU_StreamTracker_FP32 : public ::testing::Test -{ -protected: - miopen::Handle& handle = get_handle(); - miopen::StreamTracker tracker; -}; - -} // namespace - -TEST_F(GPU_StreamTracker_FP32, AcquireRelease) -{ - auto slot = tracker.acquire(handle); - ASSERT_NE(slot.stream, nullptr); - - auto saved_stream = slot.stream; - tracker.release(slot); - - auto slot2 = tracker.acquire(handle); - EXPECT_EQ(slot2.stream, saved_stream); - tracker.release(slot2); -} - -TEST_F(GPU_StreamTracker_FP32, AcquireIsNotAHandlePoolStream) -{ - // The Handle's stream-pool indices are claimed by hardcoded number elsewhere - // (MHA, RNN), so a tracker slot must never alias one of them. - constexpr int kPoolSize = 2; - handle.ReserveExtraStreamsInPool(kPoolSize); - - auto slot = tracker.acquire(handle); - ASSERT_NE(slot.stream, nullptr); - - for(int id = 0; id <= kPoolSize; ++id) - { - handle.SetStreamFromPool(id); - EXPECT_NE(slot.stream, handle.GetStream()) << "aliases stream-pool id " << id; - } - handle.SetStreamFromPool(0); - - tracker.release(slot); -} - -TEST_F(GPU_StreamTracker_FP32, AcquireGrowsPool) -{ - auto slot1 = tracker.acquire(handle); - auto slot2 = tracker.acquire(handle); - EXPECT_NE(slot1.stream, slot2.stream); - - tracker.release(slot2); - tracker.release(slot1); -} - -TEST_F(GPU_StreamTracker_FP32, AbandonAndReclaim) -{ - auto slot = tracker.acquire(handle); - - auto* dev_ptr = static_cast(nullptr); - ASSERT_EQ(hipMalloc(&dev_ptr, 64), hipSuccess); - ASSERT_EQ(hipMemsetAsync(dev_ptr, 0, 64, slot.stream), hipSuccess); - - auto abandoned_stream = slot.stream; - tracker.abandon(slot); - - ASSERT_EQ(hipStreamSynchronize(abandoned_stream), hipSuccess); - - auto reclaimed = tracker.acquire(handle); - EXPECT_EQ(reclaimed.stream, abandoned_stream); - tracker.release(reclaimed); - - ASSERT_EQ(hipFree(dev_ptr), hipSuccess); -} - -TEST_F(GPU_StreamTracker_FP32, AbandonStillDraining) -{ - auto slot = tracker.acquire(handle); - - StreamGate gate; - ASSERT_EQ(hipLaunchHostFunc(slot.stream, StreamGate::callback, &gate), hipSuccess); - - auto abandoned_stream = slot.stream; - tracker.abandon(slot); - - auto next = tracker.acquire(handle); - EXPECT_NE(next.stream, abandoned_stream); - - gate.open(); - ASSERT_EQ(hipStreamSynchronize(abandoned_stream), hipSuccess); - - // Don't release `next` yet — keep available_ empty so acquire scans draining - auto reclaimed = tracker.acquire(handle); - EXPECT_EQ(reclaimed.stream, abandoned_stream); - tracker.release(reclaimed); - tracker.release(next); -} - -TEST_F(GPU_StreamTracker_FP32, CascadeAbandonReclaim) -{ - constexpr int kCount = 4; - std::vector slots; - std::set seen; - - for(int i = 0; i < kCount; ++i) - { - auto slot = tracker.acquire(handle); - seen.insert(slot.stream); - tracker.abandon(slot); - } - - for(int i = 0; i < kCount; ++i) - { - auto slot = tracker.acquire(handle); - seen.insert(slot.stream); - slots.emplace_back(std::move(slot)); - } - - for(auto& s : slots) - tracker.release(s); - - // Every stream is idle and reclaimed, so acquiring kCount more must not - // create any stream that hasn't been handed out before. - for(int i = 0; i < kCount; ++i) - { - auto s = tracker.acquire(handle); - EXPECT_EQ(seen.count(s.stream), 1u); - tracker.release(s); - } -} - -TEST_F(GPU_StreamTracker_FP32, SweepReclaimsIdleStream) -{ - auto slot = tracker.acquire(handle); - const auto abandoned_stream = slot.stream; - tracker.abandon(std::move(slot)); - - tracker.sweep(); - - auto reclaimed = tracker.acquire(handle); - EXPECT_EQ(reclaimed.stream, abandoned_stream); - tracker.release(reclaimed); -} - -TEST_F(GPU_StreamTracker_FP32, SweepLeavesBusyStreamDraining) -{ - auto slot = tracker.acquire(handle); - - StreamGate gate; - ASSERT_EQ(hipLaunchHostFunc(slot.stream, StreamGate::callback, &gate), hipSuccess); - - auto busy_stream = slot.stream; - tracker.abandon(std::move(slot)); - - tracker.sweep(); - - // Still gated, so the slot must not have been reclaimed - auto next = tracker.acquire(handle); - EXPECT_NE(next.stream, busy_stream); - tracker.release(next); - - gate.open(); - ASSERT_EQ(hipStreamSynchronize(busy_stream), hipSuccess); - - tracker.sweep(); - - // available_ is LIFO, so the just-swept slot is on top - auto reclaimed = tracker.acquire(handle); - EXPECT_EQ(reclaimed.stream, busy_stream); - tracker.release(reclaimed); -} - -TEST_F(GPU_StreamTracker_FP32, SweepReleasesScratch) -{ - auto prev = handle.GetScratchBuffer(1); - const auto sz = (prev ? prev->size : 0) + 65536; - prev.reset(); - - auto scratch = handle.GetScratchBuffer(sz); - ASSERT_NE(scratch, nullptr); - ASSERT_EQ(scratch.use_count(), 1); - - auto slot = tracker.acquire(handle); - slot.scratch = scratch; - tracker.abandon(std::move(slot)); - ASSERT_EQ(scratch.use_count(), 2); // local + draining slot - - // No work on the stream, so sweep reclaims and drops the slot's reference - tracker.sweep(); - EXPECT_EQ(scratch.use_count(), 1); -} - -TEST_F(GPU_StreamTracker_FP32, ScratchAllocateAndReuse) -{ - auto s1 = handle.GetScratchBuffer(1024); - ASSERT_NE(s1, nullptr); - EXPECT_GE(s1->size, 1024u); - - // Same or smaller request while s1 is alive → same allocation returned - auto s2 = handle.GetScratchBuffer(s1->size); - EXPECT_EQ(s1, s2); - - auto s3 = handle.GetScratchBuffer(1); - EXPECT_EQ(s1, s3); -} - -TEST_F(GPU_StreamTracker_FP32, ScratchFreedWhenCallersRelease) -{ - // Fresh Handle so no other test can hold a ref to this scratch allocation. - miopen::Handle fresh_handle{}; - - auto scratch = fresh_handle.GetScratchBuffer(1024); - ASSERT_NE(scratch, nullptr); - EXPECT_GE(scratch->size, 1024u); - - std::weak_ptr weak = scratch; - scratch.reset(); - EXPECT_TRUE(weak.expired()); -} - -TEST_F(GPU_StreamTracker_FP32, ScratchGrows) -{ - auto s1 = handle.GetScratchBuffer(1); - ASSERT_NE(s1, nullptr); - auto* raw1 = s1->buffer.get(); - - auto s2 = handle.GetScratchBuffer(s1->size + 1); - ASSERT_NE(s2, nullptr); - EXPECT_NE(s2->buffer.get(), raw1); - EXPECT_GE(s2->size, s1->size + 1); -} - -TEST_F(GPU_StreamTracker_FP32, ScratchReturnsNullOnOversize) -{ - auto s = handle.GetScratchBuffer(handle.GetGlobalMemorySize()); - EXPECT_EQ(s, nullptr); -} - -TEST_F(GPU_StreamTracker_FP32, ScratchReturnsNullOnZero) -{ - auto s = handle.GetScratchBuffer(0); - EXPECT_EQ(s, nullptr); -} - -TEST_F(GPU_StreamTracker_FP32, ScratchSurvivesAbandon) -{ - auto prev = handle.GetScratchBuffer(1); - const auto sz = (prev ? prev->size : 0) + 65536; - prev.reset(); - - auto scratch = handle.GetScratchBuffer(sz); - ASSERT_NE(scratch, nullptr); - // Handle holds weak_ptr only; local is the sole strong ref - EXPECT_EQ(scratch.use_count(), 1); - - auto slot = tracker.acquire(handle); - slot.scratch = scratch; - EXPECT_EQ(scratch.use_count(), 2); // local + slot - - tracker.abandon(std::move(slot)); - EXPECT_EQ(scratch.use_count(), 2); // local + draining slot - - // No work on stream → hipStreamQuery succeeds → reclaim resets scratch - auto reclaimed = tracker.acquire(handle); - EXPECT_EQ(scratch.use_count(), 1); // draining slot scratch reset; only local remains - tracker.release(reclaimed); -} diff --git a/projects/miopen/test/gtest/test_categories.yaml b/projects/miopen/test/gtest/test_categories.yaml index 81515e576cc6..967d047248b3 100644 --- a/projects/miopen/test/gtest/test_categories.yaml +++ b/projects/miopen/test/gtest/test_categories.yaml @@ -64,10 +64,6 @@ positive_patterns: &positive_patterns - "*/GPU_Dropout*" - "*/GPU_MIOpenDriver*" - "*/GPU_HipGraphExist*" - # Unprefixed: these are plain TEST/TEST_F with no INSTANTIATE_TEST_SUITE_P, - # so they carry no "Prefix/" and the leading "*/" above would not match them. - - "GPU_StreamTracker_FP32*" - - "GPU_NaiveTimeout_FP32*" - "Standard/*" # include all standard tests by default diff --git a/projects/miopen/test/gtest/unit_conv_solver_naive_work_gate.cpp b/projects/miopen/test/gtest/unit_conv_solver_naive_work_gate.cpp index d96b7af98ed0..2c2dc9180e03 100644 --- a/projects/miopen/test/gtest/unit_conv_solver_naive_work_gate.cpp +++ b/projects/miopen/test/gtest/unit_conv_solver_naive_work_gate.cpp @@ -26,21 +26,36 @@ // Pure-CPU unit tests for the naive-conv work-size threshold // (ConvDirectNaiveConvExceedsWorkLimit). Naive stays *applicable* at any size -// (it is the universal fallback); this threshold is consumed at the selection -// layer (EvaluateInvokers, src/conv/solver_finders.cpp) to keep the un-tiled -// naive kernel from being *benchmarked* on shapes so large its single launch -// would trip the OS GPU watchdog (TDR) when a non-naive alternative also -// applies. See src/solver/conv/conv_direct_naive_conv.cpp for the metric. +// (it is the universal fallback); the threshold is how the three naive solvers +// answer SolverInterface::GetSpeedClass: over the limit they report +// SolverSpeedClass::ExceedsLaunchBudget, i.e. "applicable, but a single launch +// on this shape and device is estimated to outlast the OS GPU watchdog". The +// selection layers consume that answer -- FindCore/EvaluateInvokers +// (src/conv/solver_finders.cpp) for Find's mini-benchmark and +// GetSolutionsFallback (src/hip/convolutionocl.cpp) for immediate mode -- and +// benchmark only the best class present, so the un-tiled naive kernel is not +// *executed* while any better-classed alternative applies, not even a merely +// Slow one. See src/solver/conv/conv_direct_naive_conv.cpp for the metric. // -// These tests exercise the metric directly on ProblemDescriptions. They never -// construct a Handle and never launch a kernel, so they are safe on machines -// where the huge shapes would otherwise TDR. +// Two halves are tested separately because they answer different questions: +// +// * ConvDirectNaiveConvWork(problem) -- shape only. MAC count; direction- and +// layout-invariant, group-aware. +// * ConvDirectNaiveConvWorkLimit(cus, wavefront, clock_khz) -- device only. A +// wall-time budget converted to MACs via device throughput, so the limit +// scales across the CDNA/RDNA span instead of being one hand-tuned constant. +// +// These tests exercise both directly on ProblemDescriptions and synthetic device +// capabilities. They never construct a Handle and never launch a kernel, so they +// are safe on machines where the huge shapes would otherwise TDR. #include +#include #include #include +#include #include #include "unit_TensorDescriptor.hpp" @@ -54,8 +69,11 @@ using miopen::unit_tests::ConvolutionDescriptorParams; using miopen::unit_tests::TensorDescriptorParams; using Direction = miopen::conv::Direction; -// Mirror of the env override read by the gate (declared UINT64 in the library). +// Mirrors of the env overrides read by the limit model (both declared UINT64 in +// the library). MAX_WORK pins the limit outright; MAX_TIME_MS retunes the +// wall-time budget the limit is derived from. MIOPEN_LIB_ENV_VAR(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK) +MIOPEN_LIB_ENV_VAR(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS) // Build a 2D conv ProblemDescription with square spatial dims, odd square // filter, unit stride/dilation and 'same' padding (so out spatial == in @@ -114,16 +132,59 @@ MakeDepthwise2dProblem(std::size_t n, std::size_t c, std::size_t hw, std::size_t dir}; } +// Representative device capabilities, as (hw_cus, wavefront_width, clock_khz) -- +// exactly the triple the limit model consumes. Using synthetic values keeps these +// tests CPU-only while still covering the CDNA/RDNA span we have to scale across. +// hw_cus is GetMaxHardwareComputeUnits(), i.e. already CU-per-WGP doubled for gfx1*. +// +// At the 1000 ms default budget these yield roughly: +// MI300X 304 CU x wave64 x 2.1 GHz -> ~123 GMAC +// iGPU 32 WGP x wave32 x 2.9 GHz -> ~8.9 GMAC +// a ~14x spread, which is the whole point of deriving the limit instead of fixing it. +constexpr std::size_t kMi300xCus = 304, kMi300xWave = 64, kMi300xClockKhz = 2100000; +constexpr std::size_t kIgpuCus = 32, kIgpuWave = 32, kIgpuClockKhz = 2900000; + +// The fixed limit used when a device reports incomplete capabilities. Mirrors +// NAIVE_CONV_FALLBACK_MAX_WORK in the library, which is deliberately not exported: +// duplicating the value here means a change to it has to be made twice on purpose. +constexpr std::size_t kFallbackMaxWork = static_cast(16) * 1000 * 1000 * 1000; + +std::size_t WorkLimit(std::size_t cus, std::size_t wave, std::size_t clock_khz) +{ + return miopen::solver::conv::ConvDirectNaiveConvWorkLimit(cus, wave, clock_khz); +} + +std::size_t Work(const miopen::conv::ProblemDescription& p) +{ + return miopen::solver::conv::ConvDirectNaiveConvWork(p); +} + +// Stand-in for the combiner, which needs a live device. The gate is +// Work(problem) > WorkLimit(device caps); these tests drive both halves directly so +// they never construct a Handle and never launch a kernel, which is what makes them +// safe on machines where the huge shapes would otherwise TDR. +bool ExceedsWorkLimitOn(const miopen::conv::ProblemDescription& p, + std::size_t cus, + std::size_t wave, + std::size_t clock_khz) +{ + return Work(p) > WorkLimit(cus, wave, clock_khz); +} + +// Most shape-metric tests only care about the relative ordering of work counts, not +// which device is running; pin them to the small end (the iGPU) so a shape that is +// "too big" is too big for everything we ship. bool ExceedsWorkLimit(const miopen::conv::ProblemDescription& p) { - return miopen::solver::conv::ConvDirectNaiveConvExceedsWorkLimit(p); + return ExceedsWorkLimitOn(p, kIgpuCus, kIgpuWave, kIgpuClockKhz); } } // namespace // The confirmed-TDR SDXL VAE decode conv (c128 k128 768^2 3x3, N=1) is ~87 GMAC, -// far above the ~16 GMAC default limit -> over threshold (naive skipped in the -// mini-bench when a non-naive alternative also applies). +// an order of magnitude above the iGPU's ~8.9 GMAC limit -> classified +// ExceedsLaunchBudget (naive skipped in the mini-bench whenever any better-classed +// alternative also applies). TEST(CPU_ConvNaiveWorkGate_NONE, HugeShapeExceedsLimit) { EXPECT_TRUE(ExceedsWorkLimit(Make2dProblem(1, 128, 768, 128, 3, Direction::Forward))); @@ -197,3 +258,131 @@ TEST(CPU_ConvNaiveWorkGate_NONE, EnvOverrideLowersLimit) std::uint64_t{1000}); EXPECT_TRUE(ExceedsWorkLimit(small)); } + +// --------------------------------------------------------------------------- +// Device scaling. The tests above pin the device so they measure the shape half; +// these pin the shape so they measure the device half. +// --------------------------------------------------------------------------- + +// A bigger part earns a bigger budget, monotonically in every capability the model +// consumes. Stated as an ordering rather than against magic numbers so the test +// survives a recalibration of the throughput constant. +TEST(CPU_ConvNaiveWorkGate_NONE, LimitScalesWithDeviceCapability) +{ + const auto igpu = WorkLimit(kIgpuCus, kIgpuWave, kIgpuClockKhz); + const auto mi300x = WorkLimit(kMi300xCus, kMi300xWave, kMi300xClockKhz); + EXPECT_GT(mi300x, igpu); + + // Each capability independently raises the limit: double it, the limit doubles. + EXPECT_EQ(WorkLimit(2 * kIgpuCus, kIgpuWave, kIgpuClockKhz), 2 * igpu); + EXPECT_EQ(WorkLimit(kIgpuCus, 2 * kIgpuWave, kIgpuClockKhz), 2 * igpu); + EXPECT_EQ(WorkLimit(kIgpuCus, kIgpuWave, 2 * kIgpuClockKhz), 2 * igpu); +} + +// The payoff of deriving the limit instead of fixing it: the *same* shape is +// last-resort on a small part and a normal candidate on a big one. The SDXL +// VAE-decode conv (~87 GMAC) is a confirmed TDR on a laptop-class part but sits +// comfortably inside a 304-CU MI300X's ~123 GMAC second, where naive should stay in +// the running rather than being gated off a machine that can afford it. +TEST(CPU_ConvNaiveWorkGate_NONE, SameShapeSlowOnSmallDeviceNotOnLarge) +{ + const auto sdxl = Make2dProblem(1, 128, 768, 128, 3, Direction::Forward); + EXPECT_TRUE(ExceedsWorkLimitOn(sdxl, kIgpuCus, kIgpuWave, kIgpuClockKhz)); + EXPECT_FALSE(ExceedsWorkLimitOn(sdxl, kMi300xCus, kMi300xWave, kMi300xClockKhz)); +} + +// A device that cannot report a capability gives the model nothing to scale from, so +// the limit must degrade to the fixed pre-device-aware value rather than deriving a +// zero budget -- which would multiply out to 0 and gate *every* shape, including the +// tiny ones naive normally wins. +TEST(CPU_ConvNaiveWorkGate_NONE, IncompleteDeviceCapabilitiesUseFixedFallback) +{ + EXPECT_EQ(WorkLimit(0, kIgpuWave, kIgpuClockKhz), kFallbackMaxWork); + EXPECT_EQ(WorkLimit(kIgpuCus, 0, kIgpuClockKhz), kFallbackMaxWork); + EXPECT_EQ(WorkLimit(kIgpuCus, kIgpuWave, 0), kFallbackMaxWork); + + // And the fallback is a real working limit, not a disable: a small shape still + // passes under it while the huge one still trips it. + const auto small = Make2dProblem(1, 64, 56, 64, 3, Direction::Forward); + const auto huge = Make2dProblem(1, 128, 768, 128, 3, Direction::Forward); + EXPECT_LT(Work(small), kFallbackMaxWork); + EXPECT_GT(Work(huge), kFallbackMaxWork); +} + +// --------------------------------------------------------------------------- +// Time budget. The limit is a wall-time budget expressed in MACs, so the budget is +// the knob a user reaches for after changing their watchdog timeout. +// --------------------------------------------------------------------------- + +// Halving the budget halves the limit, on any device: the conversion is linear in +// time, which is what makes "set this to your TdrDelay" a usable instruction. +TEST(CPU_ConvNaiveWorkGate_NONE, TimeBudgetScalesLimitLinearly) +{ + const auto baseline = WorkLimit(kIgpuCus, kIgpuWave, kIgpuClockKhz); // 1000 ms default + { + ScopedEnvironment half(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS, + std::uint64_t{500}); + EXPECT_EQ(WorkLimit(kIgpuCus, kIgpuWave, kIgpuClockKhz), baseline / 2); + } + { + ScopedEnvironment doubled(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS, + std::uint64_t{2000}); + EXPECT_EQ(WorkLimit(kIgpuCus, kIgpuWave, kIgpuClockKhz), baseline * 2); + } +} + +// A user who raised their watchdog can raise the budget to match, putting a shape +// that was last-resort back into the running -- the same opt-in as the absolute +// MAX_WORK override but expressed in the units the watchdog actually uses. +TEST(CPU_ConvNaiveWorkGate_NONE, TimeBudgetRaisesShapeBackIntoContention) +{ + const auto sdxl = Make2dProblem(1, 128, 768, 128, 3, Direction::Forward); + ASSERT_TRUE(ExceedsWorkLimit(sdxl)); + // ~87 GMAC vs ~8.9 GMAC/s on the iGPU needs ~10 s; 30 s clears it with margin. + ScopedEnvironment budget(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS, + std::uint64_t{30000}); + EXPECT_FALSE(ExceedsWorkLimit(sdxl)); +} + +// The budget is clamped to a ceiling. This exists only so an outlandish env value +// cannot overflow the multiply into a small (or zero) limit, which would silently +// invert the override's meaning and gate everything. +TEST(CPU_ConvNaiveWorkGate_NONE, TimeBudgetIsClampedToCeiling) +{ + std::size_t at_ceiling = 0; + { + ScopedEnvironment ceiling(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS, + std::uint64_t{60} * 1000); + at_ceiling = WorkLimit(kIgpuCus, kIgpuWave, kIgpuClockKhz); + } + ScopedEnvironment absurd(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS, + std::numeric_limits::max()); + EXPECT_EQ(WorkLimit(kIgpuCus, kIgpuWave, kIgpuClockKhz), at_ceiling); +} + +// MAX_WORK wins over MAX_TIME_MS. They are not composable -- one pins the limit and +// the other retunes a derivation -- so the absolute override short-circuits first +// and the budget is never consulted. +TEST(CPU_ConvNaiveWorkGate_NONE, AbsoluteOverrideBeatsTimeBudget) +{ + ScopedEnvironment budget(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_TIME_MS, + std::uint64_t{30000}); + ScopedEnvironment absolute(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_MAX_WORK, + std::uint64_t{12345}); + EXPECT_EQ(WorkLimit(kIgpuCus, kIgpuWave, kIgpuClockKhz), 12345u); + // ...including on a device whose derived limit would be far larger. + EXPECT_EQ(WorkLimit(kMi300xCus, kMi300xWave, kMi300xClockKhz), 12345u); +} + +// The consumers of this gate express the whole last-resort policy as a comparison on +// SolverSpeedClass: benchmark the best (lowest) class that has an applicable member, +// defer everything worse. Reordering the enumerators would therefore silently invert +// the policy -- ExceedsLaunchBudget would outrank Slow, and a naive kernel this gate +// has just classified as a watchdog risk would be launched in preference to a solver +// that is merely off the pace. Pin the order so that stays a compile-visible change. +TEST(CPU_ConvNaiveWorkGate_NONE, SpeedClassOrderingPutsWatchdogRiskLast) +{ + using miopen::solver::SolverSpeedClass; + EXPECT_LT(SolverSpeedClass::Normal, SolverSpeedClass::Slow); + EXPECT_LT(SolverSpeedClass::Slow, SolverSpeedClass::ExceedsLaunchBudget); +} From 7904ee0479835b6ff6bc1695b0cc9cfbac58fbfa Mon Sep 17 00:00:00 2001 From: Brad Pepers Date: Thu, 1 Oct 2026 07:38:55 -0600 Subject: [PATCH 4/4] fix(miopen): widen NAIVE_CONV_MAX_BUDGET_MS multiply to size_t clang-tidy's bugprone-implicit-widening-of-multiplication-result flagged 60 * 1000 being evaluated in int before widening to the size_t constant. Matches the static_cast idiom already used by NAIVE_CONV_FALLBACK_MAX_WORK immediately below. No behavioural change; the product fits in int either way. --- projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp b/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp index 82bed209bd62..603736bd0093 100644 --- a/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp +++ b/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp @@ -394,7 +394,7 @@ constexpr size_t NAIVE_MAC_PER_1K_LANE_CYCLES = 3; // TdrDelay is 2 s, so 1 s leaves 2x margin. The ceiling exists only to keep an // outlandish env value from overflowing the multiply below. constexpr size_t NAIVE_CONV_DEFAULT_BUDGET_MS = 1000; -constexpr size_t NAIVE_CONV_MAX_BUDGET_MS = 60 * 1000; +constexpr size_t NAIVE_CONV_MAX_BUDGET_MS = static_cast(60) * 1000; // Limit used when the device cannot report the capabilities the model needs. This is // the fixed pre-device-aware default; it is conservative for every current part.