diff --git a/projects/miopen/CHANGELOG.md b/projects/miopen/CHANGELOG.md index b309592784e9..6999e00456a8 100644 --- a/projects/miopen/CHANGELOG.md +++ b/projects/miopen/CHANGELOG.md @@ -8,8 +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). 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 88036aaf17e1..fc3d5c532354 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -4,23 +4,27 @@ #include #include -#include +#include #include -#include +#include +#include #include #include #include #include +#include #include #include +#include #include #include #include #include #include #include +#include #include #include @@ -30,15 +34,19 @@ 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) -MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_CONV_DIRECT_MAX_SIZE, 0) + +/// 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 { @@ -211,156 +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; +bool IsDeferSlowSolversEnabled() { return !env::disabled(MIOPEN_DEBUG_DEFER_SLOW_SOLVERS); } -private: - const Handle& handle; - bool prev_profiling; -}; - -std::string NaiveSkipReason(NaiveWarmup::Status status, float best_time) +/// 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) { - 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"; + const auto it = speed_classes.find(solver_id); + return it == speed_classes.end() ? solver::SolverSpeedClass::Normal : it->second; } -} // namespace - -static NaiveWarmup TryNaiveWithTimeout(const Handle& handle, - const Invoker& invoker, - const AnyInvokeParams& invoke_ctx, - float best_time) +/// 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. +/// +/// 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. +/// +/// 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. +/// +/// 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. +/// +/// 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. +/// +/// 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) { - 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(); - } - - 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}; + 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) +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; @@ -368,39 +308,44 @@ std::vector EvaluateInvokers(const Handle& handle, if(!arch.empty()) return ret; - const auto is_naive_solver = [](const solver::ConvSolution& s) { - return s.solver_id.find("Naive") != std::string::npos; + 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); - const bool cutoff_naive = defer_naive && is_naive && non_naive_succeeded; + const auto speed_class = speed_class_of(sol); + if(ShouldSkipSlowBenchmark(speed_class, benchmark_class)) + { + MIOPEN_LOG_I("Skipping last-resort (slow) Solver: " << algorithm_name.ToString() << ":" + << sol.solver_id); + continue; + } + if(speed_class != solver::SolverSpeedClass::Normal) + { + // 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); + } if(!conv::IsEnoughWorkspace( "EvaluateInvokers", solver::Id{sol.solver_id}, sol.workspace_sz, &invoke_ctx)) @@ -469,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); @@ -551,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) { @@ -635,21 +562,85 @@ FindCoreResult FindCore(const AnyInvokeParams& invoke_ctx, ret.solutions.reserve(total); - bool non_naive_succeeded = false; - for(const auto& ss : solutions) + // 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); + if(conv_problem != nullptr && IsDeferSlowSolversEnabled()) + { + 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); + } + } + } + + // 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) { - auto evaluated = EvaluateInvokers(handle, - ss.second, - ss.first, - network_config, - invoke_ctx, - ret, - force_attach_binary, - non_naive_succeeded); - - ret.solutions.insert(ret.solutions.end(), - std::make_move_iterator(evaluated.begin()), - std::make_move_iterator(evaluated.end())); + for(const auto& ss : solutions) + { + auto evaluated = EvaluateInvokers(handle, + ss.second, + ss.first, + network_config, + invoke_ctx, + ret, + force_attach_binary, + 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; @@ -657,37 +648,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 +659,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..bbdd61b75b65 100644 --- a/projects/miopen/src/env_debug.cpp +++ b/projects/miopen/src/env_debug.cpp @@ -50,11 +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) @@ -73,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 { @@ -193,12 +194,17 @@ 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}, {MIOPEN_DEBUG_CONV_GEMM.GetName(), MIOPEN_DEBUG_CONV_GEMM}, {MIOPEN_DEBUG_CONV_IMPLICIT_GEMM.GetName(), MIOPEN_DEBUG_CONV_IMPLICIT_GEMM}, {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}, @@ -212,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 db4e2141649d..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,14 +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); +/// 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 202d92786afe..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,6 +52,32 @@ std::string ConvDirectNaiveConvCompileOption(const ExecutionContext& ctx, bool ConvDirectNaiveConvIsApplicableByKernelType(const ExecutionContext&, const miopen::conv::ProblemDescription&); +/// 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 ExecutionContext&, + 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/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 8025014dcfb1..603736bd0093 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,19 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_USE_PACKED_KERNELS); +// 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 { @@ -310,6 +323,150 @@ bool ConvDirectNaiveConvIsApplicableByKernelType(const ExecutionContext& ctx, return true; } +// 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(); + 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()); + + 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 = 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. +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 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 << " for " + << handle.GetDeviceName() + << " -- defer from selection 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/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 new file mode 100644 index 000000000000..2c2dc9180e03 --- /dev/null +++ b/projects/miopen/test/gtest/unit_conv_solver_naive_work_gate.cpp @@ -0,0 +1,388 @@ +/******************************************************************************* + * + * 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); 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. +// +// 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" +#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; + +// 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 +// 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}; +} + +// 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 ExceedsWorkLimitOn(p, kIgpuCus, kIgpuWave, kIgpuClockKhz); +} + +} // namespace + +// The confirmed-TDR SDXL VAE decode conv (c128 k128 768^2 3x3, N=1) is ~87 GMAC, +// 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))); +} + +// 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)); +} + +// --------------------------------------------------------------------------- +// 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); +}