From 372124127e3a3c486431a82befbfdb352b24bb2f Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Wed, 15 Jul 2026 13:12:00 -0400 Subject: [PATCH 01/20] feat(miopen): time-bound naive solver evaluation with disposable stream bailout Replace the hard skip of naive convolution solvers (MIOPEN_NAIVE_DISABLE_IF_ALT) with a timeout-based approach (MIOPEN_NAIVE_TIMEOUT). When a non-naive solver has already succeeded, the naive solver's warmup is launched on a disposable HIP stream with a wall-clock budget of 3x the best non-naive time. If the warmup exceeds this budget, the stream is abandoned to a detached background thread for async cleanup and the solver is skipped. If it finishes in time, normal timed benchmarking proceeds on the root stream. This preserves the find-time savings of deferring slow naive solvers while still allowing naive to win when it is genuinely the fastest option. Co-Authored-By: Claude Sonnet 4 --- projects/miopen/CHANGELOG.md | 2 +- projects/miopen/src/conv/solver_finders.cpp | 103 +++++++++++++--- projects/miopen/src/env_debug.cpp | 4 +- .../test/gtest/naive_disable_if_alt.cpp | 112 ++++++++---------- 4 files changed, 140 insertions(+), 81 deletions(-) diff --git a/projects/miopen/CHANGELOG.md b/projects/miopen/CHANGELOG.md index d0287342e9a6..78181d868c67 100644 --- a/projects/miopen/CHANGELOG.md +++ b/projects/miopen/CHANGELOG.md @@ -8,7 +8,7 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj * Added an informational log when the Composable Kernel (CK) dynamic library is loaded. ### Changed * [Conv] Load the Composable Kernel (CK) dynamic library during handle creation (`miopenCreate`) instead of on first CK use, so the one-time library-load cost is paid at a predictable setup point rather than as an unpredictable mid-run stall. -* [Conv] Naive convolution solvers are now skipped by default during find when any non-naive solver succeeds across any algorithm. Set `MIOPEN_NAIVE_DISABLE_IF_ALT=0` to restore the previous behavior. +* [Conv] Naive convolution solvers are now time-bounded during find when a non-naive solver has already succeeded. The warmup run is launched on a disposable stream with a budget of 3x the best non-naive time; if it exceeds this budget the solver is skipped. Set `MIOPEN_NAIVE_TIMEOUT=0` to disable this behavior. ### Resolved Issues * [RNN] Fix RNN workspace tensor descriptor int overflow diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index c13ad7075fe7..5f1be27ed640 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -4,9 +4,15 @@ #include #include +#include #include +#include #include +#include +#include + +#include #include #include #include @@ -26,7 +32,7 @@ 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_DISABLE_IF_ALT, true) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_NAIVE_TIMEOUT, true) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_SEARCH_CUTOFF, false) MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_FIND_SKIP_PCT, 130) MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_CONV_DIRECT_MAX_SIZE, 0) @@ -202,6 +208,63 @@ const std::vector>& GetConvSolverFinders() } // namespace conv +/// Launch a naive solver's warmup on a disposable stream with a wall-clock +/// budget of 3x the best non-naive time. Returns the warmup elapsed time in +/// ms on success, or < 0 if the kernel exceeded the budget (in which case the +/// stream is handed to a detached thread for async cleanup). +static float TryNaiveWithTimeout(const Handle& handle, + const Invoker& invoker, + const AnyInvokeParams& invoke_ctx, + float best_time) +{ + hipStream_t naive_stream = nullptr; + auto status = hipStreamCreateWithFlags(&naive_stream, hipStreamNonBlocking); + if(status != hipSuccess) + MIOPEN_THROW_HIP_STATUS(status, "Failed to create naive bailout stream"); + + const auto orig_stream = handle.GetStream(); + handle.SetStream(naive_stream); + handle.EnableProfiling(false); + + HipEventPtr ev_start = make_hip_event(); + HipEventPtr ev_stop = make_hip_event(); + hipEventRecord(ev_start.get(), naive_stream); + invoker(handle, invoke_ctx); + hipEventRecord(ev_stop.get(), naive_stream); + + const float naive_budget = best_time * 3.0f; + 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(); + } + + handle.SetStream(orig_stream); + handle.EnableProfiling(true); + + if(!finished) + { + std::thread([naive_stream]() { + hipStreamSynchronize(naive_stream); + hipStreamDestroy(naive_stream); + }).detach(); + return -1.0f; + } + + float warmup_elapsed = 0.0f; + hipEventElapsedTime(&warmup_elapsed, ev_start.get(), ev_stop.get()); + hipStreamDestroy(naive_stream); + return warmup_elapsed; +} + /// Register invoker only for the best solution within algorithm. std::vector EvaluateInvokers(const Handle& handle, const std::vector& solutions, @@ -222,11 +285,11 @@ std::vector EvaluateInvokers(const Handle& handle, return s.solver_id.find("Naive") != std::string::npos; }; - bool naive_disable = env::value(MIOPEN_NAIVE_DISABLE_IF_ALT); + 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_disable && + naive_timeout && (non_naive_succeeded || std::any_of(solutions.begin(), solutions.end(), [&](const auto& s) { return !is_naive_solver(s); })); @@ -249,18 +312,8 @@ std::vector EvaluateInvokers(const Handle& handle, { const auto& sol = solutions[idx]; - const bool is_naive = is_naive_solver(sol); - if(naive_disable && is_naive) - { - if(defer_naive && non_naive_succeeded) - { - MIOPEN_LOG_I("Skipping Naive Solver: " << algorithm_name.ToString() << ":" - << sol.solver_id); - continue; - } - MIOPEN_LOG_I("Unable to Skip Naive Solver: " << algorithm_name.ToString() << ":" - << sol.solver_id); - } + const bool is_naive = is_naive_solver(sol); + const bool cutoff_naive = defer_naive && is_naive && non_naive_succeeded; if(!conv::IsEnoughWorkspace( "EvaluateInvokers", solver::Id{sol.solver_id}, sol.workspace_sz, &invoke_ctx)) @@ -329,6 +382,22 @@ std::vector EvaluateInvokers(const Handle& handle, int i = 0; samples.clear(); + if(cutoff_naive) + { + const float warmup = + TryNaiveWithTimeout(handle, invoker, invoke_ctx, + core_result.find_search_best_time); + if(warmup < 0) + { + MIOPEN_LOG_I("Timeout: Naive Solver " + << algorithm_name.ToString() << ":" << sol.solver_id + << " exceeded 3x best budget"); + continue; + } + first_elapsed = warmup; + i = 1; + } + while(i < N_RUNS_MAX && elapsed < TIME_MS_MAX) { invoker(handle, invoke_ctx); @@ -337,8 +406,8 @@ std::vector EvaluateInvokers(const Handle& handle, if(i > 0) { samples.push_back(handle.GetKernelTime()); - if(i == 1 && using_search_cutoff && samples.front() > 1.0f && - samples.front() > skip_time) + if(i == 1 && using_search_cutoff && + samples.front() > 1.0f && samples.front() > skip_time) { MIOPEN_LOG_I("Skipping (Slow) Solver: " << algorithm_name.ToString() << ":" << sol.solver_id << " " diff --git a/projects/miopen/src/env_debug.cpp b/projects/miopen/src/env_debug.cpp index 62cad0c69e0e..71ee4a29cb15 100644 --- a/projects/miopen/src/env_debug.cpp +++ b/projects/miopen/src/env_debug.cpp @@ -71,7 +71,7 @@ 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_DISABLE_IF_ALT, true) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_NAIVE_TIMEOUT, true) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_SKIP_ASAN_DISABLED_TESTS, true) namespace miopen { @@ -208,7 +208,7 @@ 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_DISABLE_IF_ALT.GetName(), MIOPEN_NAIVE_DISABLE_IF_ALT}, + {MIOPEN_NAIVE_TIMEOUT.GetName(), MIOPEN_NAIVE_TIMEOUT}, {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/test/gtest/naive_disable_if_alt.cpp b/projects/miopen/test/gtest/naive_disable_if_alt.cpp index 2db35bd86873..dfca67bf19b0 100644 --- a/projects/miopen/test/gtest/naive_disable_if_alt.cpp +++ b/projects/miopen/test/gtest/naive_disable_if_alt.cpp @@ -1,30 +1,34 @@ // Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT -// Tests for MIOPEN_NAIVE_DISABLE_IF_ALT behavior in the find path. +// Tests for MIOPEN_NAIVE_TIMEOUT behavior in the find path. // -// Paths through EvaluateInvokers naive-deferring logic (solver_finders.cpp): +// 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. // -// Path naive_disable defer_naive non_naive_succeeded Outcome +// 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 SKIPPED (intended fast path) +// D true true true naive timed — may pass or be skipped // -// Tests cover paths A, B, and D. Path C is not exercised: -// Path C: requires fault injection into the invoker execution path at runtime. +// 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. GetEnvFindOnlySolver() re-evaluates the env var on every call -// under MIOPEN_BUILD_TESTING (no static cache) so ScopedEnvironment takes effect. +// 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. +// 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 for the -// duration of the test, ensuring TryLoad always calls the regenerator (FindCore / EvaluateInvokers) -// rather than returning a cached find DB entry. +// Each GPU test uses ScopedFindDb to disable miopen::debug::testing_find_db_enabled +// so TryLoad always calls the regenerator (FindCore / EvaluateInvokers). #include @@ -37,16 +41,10 @@ #include "gtest_common.hpp" #include "../tensor_holder.hpp" -MIOPEN_LIB_ENV_VAR(MIOPEN_NAIVE_DISABLE_IF_ALT) +MIOPEN_LIB_ENV_VAR(MIOPEN_NAIVE_TIMEOUT) namespace { -// RAII: disables the user find DB for the duration of the test so FindCore / -// EvaluateInvokers always runs rather than returning a cached result. -// miopen::debug::testing_find_db_enabled is declared in and -// checked in TryLoad: when false, construct_db() returns null → in_sync=false → -// regenerator() (FindCore) is always called. Saves and restores the previous value -// so tests that run before or after us with find DB disabled are not disturbed. struct ScopedFindDb { ScopedFindDb() : prev_(miopen::debug::testing_find_db_enabled) @@ -58,9 +56,6 @@ struct ScopedFindDb bool prev_; }; -// Returns solver names (fastest first) from miopenFindSolutions for a small NCHW FP32 forward -// convolution: 1×64×14×14 input, 64×3×3 filter, pad=1, stride=1. -// This shape attracts ImplicitGEMM/Winograd on all supported GPUs, ensuring a non-naive winner. std::vector RunFind(miopenHandle_t handle, size_t max_solutions = 8) { tensor x{1, 64, 14, 14}; @@ -104,35 +99,15 @@ bool IsNaive(const std::string& name) { return name.find("Naive") != std::string } // namespace -// With MIOPEN_NAIVE_DISABLE_IF_ALT=1, no naive solver should appear anywhere in the results -// for a shape that has non-naive alternatives. -TEST(GPU_NaiveDisableIfAlt_FP32, NaiveAbsentFromAllResultsWhenDisabled) +// 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; - ScopedFindDb no_cache; // disables find DB so FindCore / EvaluateInvokers always runs - ScopedEnvironment guard_naive(MIOPEN_NAIVE_DISABLE_IF_ALT, true); - auto solvers = RunFind(handle); - - ASSERT_FALSE(solvers.empty()) << "miopenFindSolutions returned no results"; - for(const auto& s : solvers) - { - EXPECT_FALSE(IsNaive(s)) - << "Naive solver found in results with MIOPEN_NAIVE_DISABLE_IF_ALT=1: " << s; - } -} - -// With MIOPEN_NAIVE_DISABLE_IF_ALT=0 (opt-out), a naive solver must appear somewhere in the -// results, confirming that the skip logic in EvaluateInvokers did not fire. -// This tests the opt-out code path directly without depending on which solver wins by timing. -TEST(GPU_NaiveDisableIfAlt_FP32, NaivePresentInResultsWhenAllowed) -{ - auto& handle_ref = get_handle(); - miopenHandle_t handle = &handle_ref; - - ScopedFindDb no_cache; // disables find DB so FindCore / EvaluateInvokers always runs - ScopedEnvironment guard_naive(MIOPEN_NAIVE_DISABLE_IF_ALT, false); + ScopedFindDb no_cache; + ScopedEnvironment guard_naive(MIOPEN_NAIVE_TIMEOUT, false); auto solvers = RunFind(handle); ASSERT_FALSE(solvers.empty()) << "miopenFindSolutions returned no results"; @@ -141,20 +116,20 @@ TEST(GPU_NaiveDisableIfAlt_FP32, NaivePresentInResultsWhenAllowed) 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_DISABLE_IF_ALT=0" + << "Expected at least one naive solver in results with MIOPEN_NAIVE_TIMEOUT=0" << "; got: [" << all_names << "]"; } -// With MIOPEN_NAIVE_DISABLE_IF_ALT=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 -// skip guard cannot fire (path B). Find must return a naive result rather than nothing. -TEST(GPU_NaiveDisableIfAlt_FP32, NaiveNotSuppressedWhenOnlySolver) +// 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_DISABLE_IF_ALT, true); + ScopedEnvironment guard_naive(MIOPEN_NAIVE_TIMEOUT, true); ScopedEnvironment only_naive(MIOPEN_DEBUG_FIND_ONLY_SOLVER, std::string{"ConvDirectNaiveConvFwd"}); auto solvers = RunFind(handle); @@ -167,16 +142,31 @@ TEST(GPU_NaiveDisableIfAlt_FP32, NaiveNotSuppressedWhenOnlySolver) << "Expected only naive solvers when MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvFwd"; } -// CPU-only: verify that MIOPEN_NAIVE_DISABLE_IF_ALT can be set and read back via the debug +// 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"; +} + +// 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_NaiveDisableIfAlt_NONE, EnvVarRoundTrips) +TEST(CPU_NaiveTimeout_NONE, EnvVarRoundTrips) { { - ScopedEnvironment guard(MIOPEN_NAIVE_DISABLE_IF_ALT, true); - EXPECT_TRUE(lib_env::value(MIOPEN_NAIVE_DISABLE_IF_ALT)); + ScopedEnvironment guard(MIOPEN_NAIVE_TIMEOUT, true); + EXPECT_TRUE(lib_env::value(MIOPEN_NAIVE_TIMEOUT)); } { - ScopedEnvironment guard(MIOPEN_NAIVE_DISABLE_IF_ALT, false); - EXPECT_FALSE(lib_env::value(MIOPEN_NAIVE_DISABLE_IF_ALT)); + ScopedEnvironment guard(MIOPEN_NAIVE_TIMEOUT, false); + EXPECT_FALSE(lib_env::value(MIOPEN_NAIVE_TIMEOUT)); } } From b15f103ff22a95ea0d6214eebfda5d40fde246a7 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 16 Jul 2026 14:51:13 -0400 Subject: [PATCH 02/20] feat(miopen): extract StreamTracker from solver_finders; add unit tests Move the naive timeout stream pool tracker out of solver_finders.cpp into a reusable StreamTracker class owned by the Handle. Eliminates a leaked global map keyed by Handle pointer and makes pool state naturally per-handle with no cross-thread concerns. Also registers MIOPEN_NAIVE_TIMEOUT_FACTOR in the env var system and adds five GTest cases exercising acquire/release, pool growth, abandon with async drain, and cascade reclaim. Co-Authored-By: Claude Sonnet 4 --- projects/miopen/src/conv/solver_finders.cpp | 66 +++++---- projects/miopen/src/env_debug.cpp | 2 + projects/miopen/src/hip/handlehip.cpp | 32 +++++ projects/miopen/src/include/miopen/handle.hpp | 3 + .../src/include/miopen/stream_tracker.hpp | 34 +++++ projects/miopen/test/gtest/stream_tracker.cpp | 128 ++++++++++++++++++ 6 files changed, 231 insertions(+), 34 deletions(-) create mode 100644 projects/miopen/src/include/miopen/stream_tracker.hpp create mode 100644 projects/miopen/test/gtest/stream_tracker.cpp diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index 5f1be27ed640..eb38b08cf350 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -33,6 +33,7 @@ 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_DECLARE_ENV_VAR_UINT64(MIOPEN_FIND_SKIP_PCT, 130) MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_CONV_DIRECT_MAX_SIZE, 0) @@ -208,34 +209,33 @@ const std::vector>& GetConvSolverFinders() } // namespace conv -/// Launch a naive solver's warmup on a disposable stream with a wall-clock -/// budget of 3x the best non-naive time. Returns the warmup elapsed time in -/// ms on success, or < 0 if the kernel exceeded the budget (in which case the -/// stream is handed to a detached thread for async cleanup). static float TryNaiveWithTimeout(const Handle& handle, const Invoker& invoker, const AnyInvokeParams& invoke_ctx, float best_time) { - hipStream_t naive_stream = nullptr; - auto status = hipStreamCreateWithFlags(&naive_stream, hipStreamNonBlocking); - if(status != hipSuccess) - MIOPEN_THROW_HIP_STATUS(status, "Failed to create naive bailout stream"); + auto& tracker = handle.GetStreamTracker(); + auto slot = tracker.acquire(handle); - const auto orig_stream = handle.GetStream(); - handle.SetStream(naive_stream); + handle.SetStreamFromPool(slot.pool_id); handle.EnableProfiling(false); HipEventPtr ev_start = make_hip_event(); HipEventPtr ev_stop = make_hip_event(); - hipEventRecord(ev_start.get(), naive_stream); - invoker(handle, invoke_ctx); - hipEventRecord(ev_stop.get(), naive_stream); - const float naive_budget = best_time * 3.0f; - const auto deadline = - std::chrono::steady_clock::now() + - std::chrono::microseconds(static_cast(naive_budget * 1000)); + 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, invoke_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"); + + 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) { @@ -247,22 +247,19 @@ static float TryNaiveWithTimeout(const Handle& handle, std::this_thread::yield(); } - handle.SetStream(orig_stream); + handle.SetStreamFromPool(0); handle.EnableProfiling(true); - if(!finished) + if(finished) { - std::thread([naive_stream]() { - hipStreamSynchronize(naive_stream); - hipStreamDestroy(naive_stream); - }).detach(); - return -1.0f; + tracker.release(slot); + float warmup_elapsed = 0.0f; + (void)hipEventElapsedTime(&warmup_elapsed, ev_start.get(), ev_stop.get()); + return warmup_elapsed; } - float warmup_elapsed = 0.0f; - hipEventElapsedTime(&warmup_elapsed, ev_start.get(), ev_stop.get()); - hipStreamDestroy(naive_stream); - return warmup_elapsed; + tracker.abandon(slot); + return -1.0f; } /// Register invoker only for the best solution within algorithm. @@ -384,14 +381,15 @@ std::vector EvaluateInvokers(const Handle& handle, if(cutoff_naive) { - const float warmup = - TryNaiveWithTimeout(handle, invoker, invoke_ctx, - core_result.find_search_best_time); + const float warmup = TryNaiveWithTimeout( + handle, invoker, invoke_ctx, core_result.find_search_best_time); if(warmup < 0) { MIOPEN_LOG_I("Timeout: Naive Solver " << algorithm_name.ToString() << ":" << sol.solver_id - << " exceeded 3x best budget"); + << " exceeded " << env::value(MIOPEN_NAIVE_TIMEOUT_FACTOR) + << "% best budget (" << core_result.find_search_best_time + << " ms)"); continue; } first_elapsed = warmup; @@ -406,8 +404,8 @@ std::vector EvaluateInvokers(const Handle& handle, if(i > 0) { samples.push_back(handle.GetKernelTime()); - if(i == 1 && using_search_cutoff && - samples.front() > 1.0f && samples.front() > skip_time) + if(i == 1 && using_search_cutoff && samples.front() > 1.0f && + samples.front() > skip_time) { MIOPEN_LOG_I("Skipping (Slow) Solver: " << algorithm_name.ToString() << ":" << sol.solver_id << " " diff --git a/projects/miopen/src/env_debug.cpp b/projects/miopen/src/env_debug.cpp index 71ee4a29cb15..75560a37fd1f 100644 --- a/projects/miopen/src/env_debug.cpp +++ b/projects/miopen/src/env_debug.cpp @@ -72,6 +72,7 @@ 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 { @@ -209,6 +210,7 @@ const LibEnvVar& FindEnvVariable(std::string_view name) {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/handlehip.cpp b/projects/miopen/src/hip/handlehip.cpp index 854492d37a02..79b1e2816f45 100644 --- a/projects/miopen/src/hip/handlehip.cpp +++ b/projects/miopen/src/hip/handlehip.cpp @@ -283,6 +283,7 @@ struct HandleImpl Allocator allocator{}; KernelCache cache; TargetProperties target_properties; + mutable StreamTracker stream_tracker_; }; Handle::Handle(miopenAcceleratorQueue_t stream) : impl(std::make_unique()) @@ -400,6 +401,37 @@ miopenAcceleratorQueue_t Handle::GetStream() const return this->impl->ms_resourse_ptr->stream_pool.at(meopenHandle_current_stream_id - 1).get(); } +StreamTracker& Handle::GetStreamTracker() const { return impl->stream_tracker_; } + +StreamTracker::Slot StreamTracker::acquire(const Handle& handle) +{ + if(available_.empty()) + { + for(auto it = draining_.begin(); it != draining_.end();) + { + if(hipStreamQuery(it->stream) == hipSuccess) + { + available_.push_back(*it); + it = draining_.erase(it); + } + else + ++it; + } + } + + if(!available_.empty()) + { + auto slot = available_.back(); + available_.pop_back(); + return slot; + } + + int id = next_id_++; + handle.ReserveExtraStreamsInPool(id); + handle.SetStreamFromPool(id); + return {id, handle.GetStream()}; +} + void Handle::SetAllocator(miopenAllocatorFunction allocator, miopenDeallocatorFunction deallocator, void* allocatorContext) const diff --git a/projects/miopen/src/include/miopen/handle.hpp b/projects/miopen/src/include/miopen/handle.hpp index 04b33054f6a9..d038c11f0bfc 100644 --- a/projects/miopen/src/include/miopen/handle.hpp +++ b/projects/miopen/src/include/miopen/handle.hpp @@ -38,6 +38,7 @@ #include #include #include +#include #include #include @@ -94,6 +95,8 @@ struct MIOPEN_EXPORT Handle : miopenHandle void SetStreamFromPool(int streamID) const; void ReserveExtraStreamsInPool(int cnt) const; + StreamTracker& GetStreamTracker() const; + void SetAllocator(miopenAllocatorFunction allocator, miopenDeallocatorFunction deallocator, void* allocatorContext) const; diff --git a/projects/miopen/src/include/miopen/stream_tracker.hpp b/projects/miopen/src/include/miopen/stream_tracker.hpp new file mode 100644 index 000000000000..0893ed39f5a1 --- /dev/null +++ b/projects/miopen/src/include/miopen/stream_tracker.hpp @@ -0,0 +1,34 @@ +// 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 + +namespace miopen { + +struct Handle; + +struct StreamTracker +{ + struct Slot + { + int pool_id; + hipStream_t stream; + }; + + Slot acquire(const Handle& handle); + void release(Slot slot) { available_.push_back(slot); } + void abandon(Slot slot) { draining_.push_back(slot); } + +private: + std::vector available_; + std::vector draining_; + int next_id_ = 1; +}; + +} // namespace miopen + +#endif // GUARD_MIOPEN_STREAM_TRACKER_HPP_ diff --git a/projects/miopen/test/gtest/stream_tracker.cpp b/projects/miopen/test/gtest/stream_tracker.cpp new file mode 100644 index 000000000000..efe12f74b919 --- /dev/null +++ b/projects/miopen/test/gtest/stream_tracker.cpp @@ -0,0 +1,128 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include "get_handle.hpp" + +#include +#include + +#include + +namespace { + +class GPU_StreamTracker_NONE : public ::testing::Test +{ +protected: + miopen::Handle& handle = get_handle(); + miopen::StreamTracker& tracker = handle.GetStreamTracker(); +}; + +} // namespace + +TEST_F(GPU_StreamTracker_NONE, AcquireRelease) +{ + auto slot = tracker.acquire(handle); + ASSERT_GT(slot.pool_id, 0); + ASSERT_NE(slot.stream, nullptr); + + int saved_id = slot.pool_id; + tracker.release(slot); + + auto slot2 = tracker.acquire(handle); + EXPECT_EQ(slot2.pool_id, saved_id); + tracker.release(slot2); +} + +TEST_F(GPU_StreamTracker_NONE, AcquireGrowsPool) +{ + auto slot1 = tracker.acquire(handle); + auto slot2 = tracker.acquire(handle); + EXPECT_NE(slot1.pool_id, slot2.pool_id); + + tracker.release(slot2); + tracker.release(slot1); +} + +TEST_F(GPU_StreamTracker_NONE, 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); + + int abandoned_id = slot.pool_id; + tracker.abandon(slot); + + ASSERT_EQ(hipStreamSynchronize(slot.stream), hipSuccess); + + auto reclaimed = tracker.acquire(handle); + EXPECT_EQ(reclaimed.pool_id, abandoned_id); + tracker.release(reclaimed); + + hipFree(dev_ptr); +} + +TEST_F(GPU_StreamTracker_NONE, AbandonStillDraining) +{ + auto slot = tracker.acquire(handle); + + auto* dev_ptr = static_cast(nullptr); + size_t large_size = 256 * 1024 * 1024; + ASSERT_EQ(hipMalloc(&dev_ptr, large_size), hipSuccess); + for(int i = 0; i < 64; ++i) + hipMemsetAsync(dev_ptr, 0, large_size, slot.stream); + + int abandoned_id = slot.pool_id; + tracker.abandon(slot); + + auto next = tracker.acquire(handle); + EXPECT_NE(next.pool_id, abandoned_id); + tracker.release(next); + + ASSERT_EQ(hipStreamSynchronize(slot.stream), hipSuccess); + + auto reclaimed = tracker.acquire(handle); + EXPECT_EQ(reclaimed.pool_id, abandoned_id); + tracker.release(reclaimed); + + hipFree(dev_ptr); +} + +TEST_F(GPU_StreamTracker_NONE, CascadeAbandonReclaim) +{ + constexpr int kCount = 4; + std::vector slots; + std::vector abandoned_ids; + + for(int i = 0; i < kCount; ++i) + { + auto slot = tracker.acquire(handle); + abandoned_ids.push_back(slot.pool_id); + tracker.abandon(slot); + } + + for(int i = 0; i < kCount; ++i) + { + auto& s = slots.emplace_back(tracker.acquire(handle)); + // Pool may grow since prior abandons may still be draining; + // but all should eventually reclaim once drained. + } + + // Wait for everything to drain + for(auto& s : slots) + tracker.release(s); + + // Now acquire kCount — all should come from available (no new pool growth) + auto before = tracker.acquire(handle); + int max_id = before.pool_id; + tracker.release(before); + + for(int i = 0; i < kCount; ++i) + { + auto s = tracker.acquire(handle); + EXPECT_LE(s.pool_id, max_id); + tracker.release(s); + } +} From d3f3a70494a409f084196e3c455f0659fa15b8b3 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 16 Jul 2026 15:30:38 -0400 Subject: [PATCH 03/20] fix(miopen): fix StreamTracker test compile errors and rename for PR bot MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - Check hipFree return value (nodiscard) - Cast hipMemsetAsync to void in fire-and-forget loop - Remove unused variable from emplace_back - Rename stream_tracker.cpp → stream_tracker_test.cpp to satisfy PR bot unit test naming convention Co-Authored-By: Claude Sonnet 4 --- .../{stream_tracker.cpp => stream_tracker_test.cpp} | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) rename projects/miopen/test/gtest/{stream_tracker.cpp => stream_tracker_test.cpp} (91%) diff --git a/projects/miopen/test/gtest/stream_tracker.cpp b/projects/miopen/test/gtest/stream_tracker_test.cpp similarity index 91% rename from projects/miopen/test/gtest/stream_tracker.cpp rename to projects/miopen/test/gtest/stream_tracker_test.cpp index efe12f74b919..42e4ae79ffb8 100644 --- a/projects/miopen/test/gtest/stream_tracker.cpp +++ b/projects/miopen/test/gtest/stream_tracker_test.cpp @@ -61,7 +61,7 @@ TEST_F(GPU_StreamTracker_NONE, AbandonAndReclaim) EXPECT_EQ(reclaimed.pool_id, abandoned_id); tracker.release(reclaimed); - hipFree(dev_ptr); + ASSERT_EQ(hipFree(dev_ptr), hipSuccess); } TEST_F(GPU_StreamTracker_NONE, AbandonStillDraining) @@ -72,7 +72,7 @@ TEST_F(GPU_StreamTracker_NONE, AbandonStillDraining) size_t large_size = 256 * 1024 * 1024; ASSERT_EQ(hipMalloc(&dev_ptr, large_size), hipSuccess); for(int i = 0; i < 64; ++i) - hipMemsetAsync(dev_ptr, 0, large_size, slot.stream); + (void)hipMemsetAsync(dev_ptr, 0, large_size, slot.stream); int abandoned_id = slot.pool_id; tracker.abandon(slot); @@ -87,7 +87,7 @@ TEST_F(GPU_StreamTracker_NONE, AbandonStillDraining) EXPECT_EQ(reclaimed.pool_id, abandoned_id); tracker.release(reclaimed); - hipFree(dev_ptr); + ASSERT_EQ(hipFree(dev_ptr), hipSuccess); } TEST_F(GPU_StreamTracker_NONE, CascadeAbandonReclaim) @@ -105,9 +105,7 @@ TEST_F(GPU_StreamTracker_NONE, CascadeAbandonReclaim) for(int i = 0; i < kCount; ++i) { - auto& s = slots.emplace_back(tracker.acquire(handle)); - // Pool may grow since prior abandons may still be draining; - // but all should eventually reclaim once drained. + slots.emplace_back(tracker.acquire(handle)); } // Wait for everything to drain From af4b26ea8a7aae54fee972a57b105e585a13ef65 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 16 Jul 2026 15:37:35 -0400 Subject: [PATCH 04/20] style(miopen): apply clang-format to naive_disable_if_alt test Co-Authored-By: Claude Opus 4 --- projects/miopen/test/gtest/naive_disable_if_alt.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/miopen/test/gtest/naive_disable_if_alt.cpp b/projects/miopen/test/gtest/naive_disable_if_alt.cpp index dfca67bf19b0..b9841bb563b7 100644 --- a/projects/miopen/test/gtest/naive_disable_if_alt.cpp +++ b/projects/miopen/test/gtest/naive_disable_if_alt.cpp @@ -116,8 +116,8 @@ TEST(GPU_NaiveTimeout_FP32, NaivePresentWhenTimeoutDisabled) 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 << "]"; + << "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 From 2e0e5482f636300bcaf84c9f25fa504ebc791520 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 16 Jul 2026 15:58:45 -0400 Subject: [PATCH 05/20] fix(miopen): export StreamTracker symbol for test linkage Co-Authored-By: Claude Opus 4 --- projects/miopen/src/include/miopen/stream_tracker.hpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/projects/miopen/src/include/miopen/stream_tracker.hpp b/projects/miopen/src/include/miopen/stream_tracker.hpp index 0893ed39f5a1..2a87c24f9671 100644 --- a/projects/miopen/src/include/miopen/stream_tracker.hpp +++ b/projects/miopen/src/include/miopen/stream_tracker.hpp @@ -4,6 +4,8 @@ #ifndef GUARD_MIOPEN_STREAM_TRACKER_HPP_ #define GUARD_MIOPEN_STREAM_TRACKER_HPP_ +#include + #include #include @@ -11,7 +13,7 @@ namespace miopen { struct Handle; -struct StreamTracker +struct MIOPEN_EXPORT StreamTracker { struct Slot { From 2a2c59d83ffa13fa9b884a2c82d5dd9d609da1ed Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Thu, 16 Jul 2026 16:34:22 -0400 Subject: [PATCH 06/20] fix(miopen): rename GPU_StreamTracker_NONE to GPU_StreamTracker_FP32 GPU+NONE is rejected by check_names.py (line 98). GPU tests require a real datatype suffix. Co-Authored-By: Claude Opus 4 --- projects/miopen/test/gtest/stream_tracker_test.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/projects/miopen/test/gtest/stream_tracker_test.cpp b/projects/miopen/test/gtest/stream_tracker_test.cpp index 42e4ae79ffb8..6a799e6efd01 100644 --- a/projects/miopen/test/gtest/stream_tracker_test.cpp +++ b/projects/miopen/test/gtest/stream_tracker_test.cpp @@ -11,7 +11,7 @@ namespace { -class GPU_StreamTracker_NONE : public ::testing::Test +class GPU_StreamTracker_FP32 : public ::testing::Test { protected: miopen::Handle& handle = get_handle(); @@ -20,7 +20,7 @@ class GPU_StreamTracker_NONE : public ::testing::Test } // namespace -TEST_F(GPU_StreamTracker_NONE, AcquireRelease) +TEST_F(GPU_StreamTracker_FP32, AcquireRelease) { auto slot = tracker.acquire(handle); ASSERT_GT(slot.pool_id, 0); @@ -34,7 +34,7 @@ TEST_F(GPU_StreamTracker_NONE, AcquireRelease) tracker.release(slot2); } -TEST_F(GPU_StreamTracker_NONE, AcquireGrowsPool) +TEST_F(GPU_StreamTracker_FP32, AcquireGrowsPool) { auto slot1 = tracker.acquire(handle); auto slot2 = tracker.acquire(handle); @@ -44,7 +44,7 @@ TEST_F(GPU_StreamTracker_NONE, AcquireGrowsPool) tracker.release(slot1); } -TEST_F(GPU_StreamTracker_NONE, AbandonAndReclaim) +TEST_F(GPU_StreamTracker_FP32, AbandonAndReclaim) { auto slot = tracker.acquire(handle); @@ -64,7 +64,7 @@ TEST_F(GPU_StreamTracker_NONE, AbandonAndReclaim) ASSERT_EQ(hipFree(dev_ptr), hipSuccess); } -TEST_F(GPU_StreamTracker_NONE, AbandonStillDraining) +TEST_F(GPU_StreamTracker_FP32, AbandonStillDraining) { auto slot = tracker.acquire(handle); @@ -90,7 +90,7 @@ TEST_F(GPU_StreamTracker_NONE, AbandonStillDraining) ASSERT_EQ(hipFree(dev_ptr), hipSuccess); } -TEST_F(GPU_StreamTracker_NONE, CascadeAbandonReclaim) +TEST_F(GPU_StreamTracker_FP32, CascadeAbandonReclaim) { constexpr int kCount = 4; std::vector slots; From 547f5baef1ca1037ba317fb47835987e55db7274 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Mon, 27 Jul 2026 10:01:21 -0400 Subject: [PATCH 07/20] docs(miopen): document MIOPEN_NAIVE_TIMEOUT and MIOPEN_NAIVE_TIMEOUT_FACTOR env vars Replace removed MIOPEN_NAIVE_DISABLE_IF_ALT with the two new env vars and document their interaction as requested by reviewer. Co-Authored-By: Claude Opus 4 --- .../miopen/docs/reference/env_variables.rst | 22 +++++++++++++++---- 1 file changed, 18 insertions(+), 4 deletions(-) diff --git a/projects/miopen/docs/reference/env_variables.rst b/projects/miopen/docs/reference/env_variables.rst index 4f7cec244130..0e9b57314803 100644 --- a/projects/miopen/docs/reference/env_variables.rst +++ b/projects/miopen/docs/reference/env_variables.rst @@ -104,13 +104,27 @@ and :doc:`Performance database <../conceptual/perfdb>`. - | 1: Enable | 0 or unset: Disable - * - | ``MIOPEN_NAIVE_DISABLE_IF_ALT`` - | Skips naive convolution solvers during find when at least one - | non-naive solver succeeds across any algorithm. Naive solvers - | are still used as a fallback when no non-naive solver succeeds. + * - | ``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. - | 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_DISABLE_FIND_DB`` | Disables FindDb functionality. - | 1: Disable FindDb From c8bcd4197265bec8998daeaf9d4544eb5af640b1 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 28 Jul 2026 11:59:04 -0400 Subject: [PATCH 08/20] fix(miopen): make StreamTracker tests deterministic across GPU speeds - Use per-test StreamTracker instances to avoid cross-test state leakage - Replace racey large-memset with hipLaunchHostFunc barrier to deterministically hold the stream busy in AbandonStillDraining - Defer releasing available slot so acquire is forced to scan draining Verified: 5/5 tests pass on gfx942 (MI300X). Co-Authored-By: Claude Opus 4 --- .../miopen/test/gtest/stream_tracker_test.cpp | 43 ++++++++++++++----- 1 file changed, 33 insertions(+), 10 deletions(-) diff --git a/projects/miopen/test/gtest/stream_tracker_test.cpp b/projects/miopen/test/gtest/stream_tracker_test.cpp index 6a799e6efd01..36253dd37d2d 100644 --- a/projects/miopen/test/gtest/stream_tracker_test.cpp +++ b/projects/miopen/test/gtest/stream_tracker_test.cpp @@ -9,13 +9,39 @@ #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 = handle.GetStreamTracker(); + miopen::Handle& handle = get_handle(); + miopen::StreamTracker tracker; }; } // namespace @@ -68,26 +94,23 @@ TEST_F(GPU_StreamTracker_FP32, AbandonStillDraining) { auto slot = tracker.acquire(handle); - auto* dev_ptr = static_cast(nullptr); - size_t large_size = 256 * 1024 * 1024; - ASSERT_EQ(hipMalloc(&dev_ptr, large_size), hipSuccess); - for(int i = 0; i < 64; ++i) - (void)hipMemsetAsync(dev_ptr, 0, large_size, slot.stream); + StreamGate gate; + ASSERT_EQ(hipLaunchHostFunc(slot.stream, StreamGate::callback, &gate), hipSuccess); int abandoned_id = slot.pool_id; tracker.abandon(slot); auto next = tracker.acquire(handle); EXPECT_NE(next.pool_id, abandoned_id); - tracker.release(next); + gate.open(); ASSERT_EQ(hipStreamSynchronize(slot.stream), hipSuccess); + // Don't release `next` yet — keep available_ empty so acquire scans draining auto reclaimed = tracker.acquire(handle); EXPECT_EQ(reclaimed.pool_id, abandoned_id); tracker.release(reclaimed); - - ASSERT_EQ(hipFree(dev_ptr), hipSuccess); + tracker.release(next); } TEST_F(GPU_StreamTracker_FP32, CascadeAbandonReclaim) From 9cdaaf1441e59a8ffbb16a3ca3245ff7cac31a86 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Mon, 3 Aug 2026 15:57:52 -0400 Subject: [PATCH 09/20] fix(miopen): add scratch buffer to prevent buffer aliasing in TryNaiveWithTimeout MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit When a naive kernel times out and its stream is abandoned, the kernel continues writing to the shared output buffer while the next solver reuses those same buffers on stream 0. Redirect naive output to a per-Handle scratch buffer (capped at 1/8 device memory) so abandoned writes are harmless. Scratch lifetime is caller-driven via shared_ptr attached to StreamTracker slots — RAII frees when the last user finishes. Also adds RAII exception safety for stream/profiling restore, renames test files to match convention, and moves CHANGELOG entry to ROCm 10.1. Co-Authored-By: Claude Sonnet 4 --- projects/miopen/CHANGELOG.md | 6 +- projects/miopen/src/conv/solver_finders.cpp | 58 +++++++++++++++--- projects/miopen/src/hip/handlehip.cpp | 30 +++++++++- projects/miopen/src/include/miopen/handle.hpp | 1 + .../src/include/miopen/stream_tracker.hpp | 19 +++++- ...e_disable_if_alt.cpp => naive_timeout.cpp} | 22 +++++++ ...am_tracker_test.cpp => stream_tracker.cpp} | 60 +++++++++++++++++++ 7 files changed, 181 insertions(+), 15 deletions(-) rename projects/miopen/test/gtest/{naive_disable_if_alt.cpp => naive_timeout.cpp} (87%) rename projects/miopen/test/gtest/{stream_tracker_test.cpp => stream_tracker.cpp} (68%) diff --git a/projects/miopen/CHANGELOG.md b/projects/miopen/CHANGELOG.md index 28ac2f5baad6..a2000d30dfe2 100644 --- a/projects/miopen/CHANGELOG.md +++ b/projects/miopen/CHANGELOG.md @@ -3,6 +3,11 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/projects/MIOpen/en/latest/) +## MIOpen for ROCm 10.1 + +### Changed +* [Conv] Naive convolution solvers are now time-bounded during find when a non-naive solver has already succeeded. The warmup run is launched on a separate stream with a budget of 3x the best non-naive time; if it exceeds this budget the solver is skipped. Set `MIOPEN_NAIVE_TIMEOUT=0` to disable this behavior. + ## MIOpen 3.5.2 for ROCm 7.14.0 ### Added * Added an informational log when the Composable Kernel (CK) dynamic library is loaded. @@ -12,7 +17,6 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj ### Changed * [Conv] Load the Composable Kernel (CK) dynamic library during handle creation (`miopenCreate`) instead of on first CK use, so the one-time library-load cost is paid at a predictable setup point rather than as an unpredictable mid-run stall. -* [Conv] Naive convolution solvers are now time-bounded during find when a non-naive solver has already succeeded. The warmup run is launched on a disposable stream with a budget of 3x the best non-naive time; if it exceeds this budget the solver is skipped. Set `MIOPEN_NAIVE_TIMEOUT=0` to disable this behavior. * [Conv] Refreshed the gfx942 SystemDB with updated tuned entries from recent perf-eval runs. * [Conv] Refined the 2D AI heuristics (Tunanet solver selection and kernel-tuning-network candidate ranking) with feature-engineered inputs for gfx942 and gfx950, improving selection for grouped implicit-GEMM convolutions. diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index eb38b08cf350..efa5b9c3cb0d 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -18,7 +18,9 @@ #include #include #include +#include #include +#include #include #include @@ -214,22 +216,61 @@ static float TryNaiveWithTimeout(const Handle& handle, const AnyInvokeParams& invoke_ctx, float best_time) { + std::shared_ptr scratch; + AnyInvokeParams scratch_ctx; + + if(invoke_ctx.IsOfType()) + { + auto params = invoke_ctx.CastTo(); + scratch = handle.GetScratchBuffer(params.tensors.outDesc.GetNumBytes()); + if(!scratch) + return -1.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 -1.0f; + params.tensors.dw = scratch->buffer.get(); + scratch_ctx = AnyInvokeParams{params}; + } + else + return -1.0f; + auto& tracker = handle.GetStreamTracker(); auto slot = tracker.acquire(handle); + slot.scratch = scratch; handle.SetStreamFromPool(slot.pool_id); handle.EnableProfiling(false); + auto guard = [&] { + handle.SetStreamFromPool(0); + handle.EnableProfiling(true); + }; + HipEventPtr ev_start = make_hip_event(); HipEventPtr ev_stop = make_hip_event(); - 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, invoke_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"); + 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(...) + { + guard(); + tracker.abandon(slot); + throw; + } const float timeout_factor = static_cast(env::value(MIOPEN_NAIVE_TIMEOUT_FACTOR)) / 100.0f; @@ -247,8 +288,7 @@ static float TryNaiveWithTimeout(const Handle& handle, std::this_thread::yield(); } - handle.SetStreamFromPool(0); - handle.EnableProfiling(true); + guard(); if(finished) { diff --git a/projects/miopen/src/hip/handlehip.cpp b/projects/miopen/src/hip/handlehip.cpp index 79b1e2816f45..3e6c25dbea46 100644 --- a/projects/miopen/src/hip/handlehip.cpp +++ b/projects/miopen/src/hip/handlehip.cpp @@ -284,6 +284,8 @@ struct HandleImpl KernelCache cache; TargetProperties target_properties; mutable StreamTracker stream_tracker_; + std::shared_ptr active_scratch_; + std::size_t scratch_cap_ = 0; }; Handle::Handle(miopenAcceleratorQueue_t stream) : impl(std::make_unique()) @@ -403,6 +405,27 @@ miopenAcceleratorQueue_t Handle::GetStream() const 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(impl->active_scratch_ && sz <= impl->active_scratch_->size) + return impl->active_scratch_; + + auto alloc = std::make_shared(); + alloc->buffer = impl->allocator(sz); + alloc->size = sz; + impl->active_scratch_ = alloc; + return alloc; +} + StreamTracker::Slot StreamTracker::acquire(const Handle& handle) { if(available_.empty()) @@ -411,7 +434,8 @@ StreamTracker::Slot StreamTracker::acquire(const Handle& handle) { if(hipStreamQuery(it->stream) == hipSuccess) { - available_.push_back(*it); + it->scratch.reset(); + available_.push_back(std::move(*it)); it = draining_.erase(it); } else @@ -421,7 +445,7 @@ StreamTracker::Slot StreamTracker::acquire(const Handle& handle) if(!available_.empty()) { - auto slot = available_.back(); + auto slot = std::move(available_.back()); available_.pop_back(); return slot; } @@ -429,7 +453,7 @@ StreamTracker::Slot StreamTracker::acquire(const Handle& handle) int id = next_id_++; handle.ReserveExtraStreamsInPool(id); handle.SetStreamFromPool(id); - return {id, handle.GetStream()}; + return {id, handle.GetStream(), {}}; } void Handle::SetAllocator(miopenAllocatorFunction allocator, diff --git a/projects/miopen/src/include/miopen/handle.hpp b/projects/miopen/src/include/miopen/handle.hpp index d038c11f0bfc..2500b380ddc6 100644 --- a/projects/miopen/src/include/miopen/handle.hpp +++ b/projects/miopen/src/include/miopen/handle.hpp @@ -185,6 +185,7 @@ 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/stream_tracker.hpp b/projects/miopen/src/include/miopen/stream_tracker.hpp index 2a87c24f9671..56d38a64c22f 100644 --- a/projects/miopen/src/include/miopen/stream_tracker.hpp +++ b/projects/miopen/src/include/miopen/stream_tracker.hpp @@ -5,7 +5,9 @@ #define GUARD_MIOPEN_STREAM_TRACKER_HPP_ #include +#include +#include #include #include @@ -13,17 +15,30 @@ namespace miopen { struct Handle; +struct ScratchAllocation +{ + Allocator::ManageDataPtr buffer; + std::size_t size = 0; +}; + struct MIOPEN_EXPORT StreamTracker { struct Slot { int pool_id; hipStream_t stream; + std::shared_ptr scratch; }; Slot acquire(const Handle& handle); - void release(Slot slot) { available_.push_back(slot); } - void abandon(Slot slot) { draining_.push_back(slot); } + + void release(Slot slot) + { + slot.scratch.reset(); + available_.push_back(std::move(slot)); + } + + void abandon(Slot slot) { draining_.push_back(std::move(slot)); } private: std::vector available_; diff --git a/projects/miopen/test/gtest/naive_disable_if_alt.cpp b/projects/miopen/test/gtest/naive_timeout.cpp similarity index 87% rename from projects/miopen/test/gtest/naive_disable_if_alt.cpp rename to projects/miopen/test/gtest/naive_timeout.cpp index b9841bb563b7..a61670e8772a 100644 --- a/projects/miopen/test/gtest/naive_disable_if_alt.cpp +++ b/projects/miopen/test/gtest/naive_timeout.cpp @@ -42,6 +42,7 @@ #include "../tensor_holder.hpp" MIOPEN_LIB_ENV_VAR(MIOPEN_NAIVE_TIMEOUT) +MIOPEN_LIB_ENV_VAR(MIOPEN_NAIVE_TIMEOUT_FACTOR) namespace { @@ -157,6 +158,27 @@ TEST(GPU_NaiveTimeout_FP32, FindSucceedsWithTimeoutEnabled) 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; + + 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) diff --git a/projects/miopen/test/gtest/stream_tracker_test.cpp b/projects/miopen/test/gtest/stream_tracker.cpp similarity index 68% rename from projects/miopen/test/gtest/stream_tracker_test.cpp rename to projects/miopen/test/gtest/stream_tracker.cpp index 36253dd37d2d..71ef2125626b 100644 --- a/projects/miopen/test/gtest/stream_tracker_test.cpp +++ b/projects/miopen/test/gtest/stream_tracker.cpp @@ -147,3 +147,63 @@ TEST_F(GPU_StreamTracker_FP32, CascadeAbandonReclaim) tracker.release(s); } } + +TEST_F(GPU_StreamTracker_FP32, ScratchAllocateAndReuse) +{ + auto s1 = handle.GetScratchBuffer(1024); + ASSERT_NE(s1, nullptr); + EXPECT_GE(s1->size, 1024u); + + auto s2 = handle.GetScratchBuffer(s1->size); + EXPECT_EQ(s1, s2); + + auto s3 = handle.GetScratchBuffer(1); + EXPECT_EQ(s1, s3); +} + +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); + EXPECT_EQ(scratch.use_count(), 2); // local + handle active + + auto slot = tracker.acquire(handle); + slot.scratch = scratch; + EXPECT_EQ(scratch.use_count(), 3); // local + handle + slot + + tracker.abandon(std::move(slot)); + EXPECT_EQ(scratch.use_count(), 3); // local + handle + draining + + // No work on stream → hipStreamQuery succeeds → reclaim resets scratch + auto reclaimed = tracker.acquire(handle); + EXPECT_EQ(scratch.use_count(), 2); // local + handle + tracker.release(reclaimed); +} From 4b58cdba6aa58c10286e377f704c99a5052282aa Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Mon, 3 Aug 2026 16:04:29 -0400 Subject: [PATCH 10/20] style(miopen): apply clang-format to GetScratchBuffer Co-Authored-By: Claude Sonnet 4 --- projects/miopen/src/hip/handlehip.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/miopen/src/hip/handlehip.cpp b/projects/miopen/src/hip/handlehip.cpp index 3e6c25dbea46..78d40b70ec16 100644 --- a/projects/miopen/src/hip/handlehip.cpp +++ b/projects/miopen/src/hip/handlehip.cpp @@ -419,9 +419,9 @@ std::shared_ptr Handle::GetScratchBuffer(std::size_t sz) cons if(impl->active_scratch_ && sz <= impl->active_scratch_->size) return impl->active_scratch_; - auto alloc = std::make_shared(); - alloc->buffer = impl->allocator(sz); - alloc->size = sz; + auto alloc = std::make_shared(); + alloc->buffer = impl->allocator(sz); + alloc->size = sz; impl->active_scratch_ = alloc; return alloc; } From 8c22ff23c87e8b23a6fdadf6b0ff6e71d701a288 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 4 Aug 2026 00:41:46 -0400 Subject: [PATCH 11/20] fix(miopen): use weak_ptr for scratch buffer to avoid memory pressure during inference The Handle held a shared_ptr to the scratch buffer used by TryNaiveWithTimeout, keeping the allocation alive after the find phase completed. This reduced available device memory during steady-state inference, causing ~4-6% throughput degradation. Change to weak_ptr so the scratch buffer is freed automatically when the last caller (TryNaiveWithTimeout or a draining StreamTracker slot) releases its reference. Co-Authored-By: Claude Sonnet 4 --- projects/miopen/src/hip/handlehip.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/miopen/src/hip/handlehip.cpp b/projects/miopen/src/hip/handlehip.cpp index 78d40b70ec16..68045e3d14e4 100644 --- a/projects/miopen/src/hip/handlehip.cpp +++ b/projects/miopen/src/hip/handlehip.cpp @@ -284,7 +284,7 @@ struct HandleImpl KernelCache cache; TargetProperties target_properties; mutable StreamTracker stream_tracker_; - std::shared_ptr active_scratch_; + std::weak_ptr active_scratch_; std::size_t scratch_cap_ = 0; }; @@ -416,8 +416,8 @@ std::shared_ptr Handle::GetScratchBuffer(std::size_t sz) cons if(sz > impl->scratch_cap_) return nullptr; - if(impl->active_scratch_ && sz <= impl->active_scratch_->size) - return impl->active_scratch_; + if(auto existing = impl->active_scratch_.lock(); existing && sz <= existing->size) + return existing; auto alloc = std::make_shared(); alloc->buffer = impl->allocator(sz); From 86819a4d72a50cf293785d40c422b324871c4d7a Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 4 Aug 2026 10:50:37 -0400 Subject: [PATCH 12/20] test(miopen): update StreamTracker tests for weak_ptr scratch semantics Fix ScratchSurvivesAbandon use_count assertions to match the weak_ptr ownership model (Handle holds weak_ptr, not shared_ptr, so the count is 1 lower at each checkpoint). Add ScratchFreedWhenCallersRelease to directly verify the core fix: scratch is freed when all callers release, because the Handle's weak_ptr does not extend its lifetime. Co-Authored-By: Claude Sonnet 4 --- projects/miopen/test/gtest/stream_tracker.cpp | 28 ++++++++++++++++--- 1 file changed, 24 insertions(+), 4 deletions(-) diff --git a/projects/miopen/test/gtest/stream_tracker.cpp b/projects/miopen/test/gtest/stream_tracker.cpp index 71ef2125626b..ad8a3edfaffc 100644 --- a/projects/miopen/test/gtest/stream_tracker.cpp +++ b/projects/miopen/test/gtest/stream_tracker.cpp @@ -154,6 +154,7 @@ TEST_F(GPU_StreamTracker_FP32, ScratchAllocateAndReuse) 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); @@ -161,6 +162,24 @@ TEST_F(GPU_StreamTracker_FP32, ScratchAllocateAndReuse) EXPECT_EQ(s1, s3); } +TEST_F(GPU_StreamTracker_FP32, ScratchFreedWhenCallersRelease) +{ + // Core memory-pressure fix: Handle holds only weak_ptr, so scratch is freed + // as soon as all callers (TryNaiveWithTimeout stack frame + draining slots) drop refs. + auto scratch = handle.GetScratchBuffer(1024); + ASSERT_NE(scratch, nullptr); + EXPECT_EQ(scratch.use_count(), 1); // only local; Handle has weak_ptr + + std::weak_ptr weak = scratch; + scratch.reset(); // simulate caller (find phase) completing + EXPECT_TRUE(weak.expired()); // no strong refs remain → allocation freed + + // Next call must allocate fresh (weak_ptr expired) + auto s2 = handle.GetScratchBuffer(1024); + ASSERT_NE(s2, nullptr); + EXPECT_EQ(s2.use_count(), 1); +} + TEST_F(GPU_StreamTracker_FP32, ScratchGrows) { auto s1 = handle.GetScratchBuffer(1); @@ -193,17 +212,18 @@ TEST_F(GPU_StreamTracker_FP32, ScratchSurvivesAbandon) auto scratch = handle.GetScratchBuffer(sz); ASSERT_NE(scratch, nullptr); - EXPECT_EQ(scratch.use_count(), 2); // local + handle active + // 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(), 3); // local + handle + slot + EXPECT_EQ(scratch.use_count(), 2); // local + slot tracker.abandon(std::move(slot)); - EXPECT_EQ(scratch.use_count(), 3); // local + handle + draining + 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(), 2); // local + handle + EXPECT_EQ(scratch.use_count(), 1); // draining slot scratch reset; only local remains tracker.release(reclaimed); } From 32bc285dabb34f24be763ca6bd7bbbd3ec47f5dd Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 4 Aug 2026 10:59:47 -0400 Subject: [PATCH 13/20] style(miopen): apply clang-format to stream_tracker tests Co-Authored-By: Claude Sonnet 4 --- projects/miopen/test/gtest/stream_tracker.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/miopen/test/gtest/stream_tracker.cpp b/projects/miopen/test/gtest/stream_tracker.cpp index ad8a3edfaffc..cf0b1f7e7b95 100644 --- a/projects/miopen/test/gtest/stream_tracker.cpp +++ b/projects/miopen/test/gtest/stream_tracker.cpp @@ -171,7 +171,7 @@ TEST_F(GPU_StreamTracker_FP32, ScratchFreedWhenCallersRelease) EXPECT_EQ(scratch.use_count(), 1); // only local; Handle has weak_ptr std::weak_ptr weak = scratch; - scratch.reset(); // simulate caller (find phase) completing + scratch.reset(); // simulate caller (find phase) completing EXPECT_TRUE(weak.expired()); // no strong refs remain → allocation freed // Next call must allocate fresh (weak_ptr expired) From 82bc666dd717bbf7aaaf5109ee066b9ef2cd589f Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 4 Aug 2026 13:02:04 -0400 Subject: [PATCH 14/20] fix(miopen): qualify ScratchAllocation namespace in stream_tracker test ScratchAllocation lives in namespace miopen; the anonymous namespace in the test file requires the fully qualified name. Co-Authored-By: Claude Sonnet 4 --- projects/miopen/test/gtest/stream_tracker.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/miopen/test/gtest/stream_tracker.cpp b/projects/miopen/test/gtest/stream_tracker.cpp index cf0b1f7e7b95..da1743e454c9 100644 --- a/projects/miopen/test/gtest/stream_tracker.cpp +++ b/projects/miopen/test/gtest/stream_tracker.cpp @@ -170,7 +170,7 @@ TEST_F(GPU_StreamTracker_FP32, ScratchFreedWhenCallersRelease) ASSERT_NE(scratch, nullptr); EXPECT_EQ(scratch.use_count(), 1); // only local; Handle has weak_ptr - std::weak_ptr weak = scratch; + std::weak_ptr weak = scratch; scratch.reset(); // simulate caller (find phase) completing EXPECT_TRUE(weak.expired()); // no strong refs remain → allocation freed From dbf31b9841490b210c6965138a95c38b0f7f588a Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 4 Aug 2026 16:53:07 -0400 Subject: [PATCH 15/20] fix(miopen): make ScratchFreedWhenCallersRelease robust to shared Handle The test runs in miopen_gtest where the Handle is a process-wide singleton. Other tests in the same shard may have called GetScratchBuffer earlier, leaving a live shared_ptr that the weak_ptr can lock onto. Fix: force a fresh allocation larger than any cached scratch, then verify the weak_ptr expires once the caller drops its ref. Co-Authored-By: Claude Sonnet 4 --- projects/miopen/test/gtest/stream_tracker.cpp | 19 ++++++++----------- 1 file changed, 8 insertions(+), 11 deletions(-) diff --git a/projects/miopen/test/gtest/stream_tracker.cpp b/projects/miopen/test/gtest/stream_tracker.cpp index da1743e454c9..3e2d1b45408b 100644 --- a/projects/miopen/test/gtest/stream_tracker.cpp +++ b/projects/miopen/test/gtest/stream_tracker.cpp @@ -164,20 +164,17 @@ TEST_F(GPU_StreamTracker_FP32, ScratchAllocateAndReuse) TEST_F(GPU_StreamTracker_FP32, ScratchFreedWhenCallersRelease) { - // Core memory-pressure fix: Handle holds only weak_ptr, so scratch is freed - // as soon as all callers (TryNaiveWithTimeout stack frame + draining slots) drop refs. - auto scratch = handle.GetScratchBuffer(1024); + // Force a fresh allocation larger than anything the shared Handle has cached. + auto prior = handle.GetScratchBuffer(1); + const auto sz = (prior ? prior->size : 0) + 131072; + prior.reset(); + + auto scratch = handle.GetScratchBuffer(sz); ASSERT_NE(scratch, nullptr); - EXPECT_EQ(scratch.use_count(), 1); // only local; Handle has weak_ptr std::weak_ptr weak = scratch; - scratch.reset(); // simulate caller (find phase) completing - EXPECT_TRUE(weak.expired()); // no strong refs remain → allocation freed - - // Next call must allocate fresh (weak_ptr expired) - auto s2 = handle.GetScratchBuffer(1024); - ASSERT_NE(s2, nullptr); - EXPECT_EQ(s2.use_count(), 1); + scratch.reset(); + EXPECT_TRUE(weak.expired()); } TEST_F(GPU_StreamTracker_FP32, ScratchGrows) From 3dad1fba1464d6cd2c29fe30d34df28d0b215354 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 4 Aug 2026 17:20:57 -0400 Subject: [PATCH 16/20] fix(miopen): use fresh Handle to isolate ScratchFreedWhenCallersRelease test Use a local Handle instead of the shared singleton so no other test can hold a strong ref to the scratch allocation, making the expiry check deterministic in both standalone and monolithic test binaries. Co-Authored-By: Claude Sonnet 4 --- projects/miopen/test/gtest/stream_tracker.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/projects/miopen/test/gtest/stream_tracker.cpp b/projects/miopen/test/gtest/stream_tracker.cpp index 3e2d1b45408b..5ec278505bfd 100644 --- a/projects/miopen/test/gtest/stream_tracker.cpp +++ b/projects/miopen/test/gtest/stream_tracker.cpp @@ -164,13 +164,12 @@ TEST_F(GPU_StreamTracker_FP32, ScratchAllocateAndReuse) TEST_F(GPU_StreamTracker_FP32, ScratchFreedWhenCallersRelease) { - // Force a fresh allocation larger than anything the shared Handle has cached. - auto prior = handle.GetScratchBuffer(1); - const auto sz = (prior ? prior->size : 0) + 131072; - prior.reset(); + // Fresh Handle so no other test can hold a ref to this scratch allocation. + miopen::Handle fresh_handle{}; - auto scratch = handle.GetScratchBuffer(sz); + auto scratch = fresh_handle.GetScratchBuffer(1024); ASSERT_NE(scratch, nullptr); + EXPECT_GE(scratch->size, 1024u); std::weak_ptr weak = scratch; scratch.reset(); From b619bcf47b96c0739782fb3aa0b4b166bc784397 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Tue, 11 Aug 2026 12:58:51 -0400 Subject: [PATCH 17/20] fix(miopen): address review feedback on naive solver timeout Reclaim abandoned streams eagerly rather than only on the next acquire: sweep() is split out of acquire(), runs unconditionally there, and also runs from Handle::Run(), so scratch pinned by a timed-out evaluation is released on the next kernel launch after its stream goes idle. Block on draining streams in ~StreamTracker(). Abandoned kernels execute from code objects the Handle unloads via hipModuleUnload during teardown, so they must finish before the tracker's owner goes away. HandleImpl declares stream_tracker_ after the stream pool, so the pooled streams are still valid when the tracker destructs. The wait is logged before and after since it can be long. Replace the -1.0f sentinel from TryNaiveWithTimeout with an explicit status enum, so a genuine timeout is no longer indistinguishable from an oversized scratch request or an unsupported invoke-params type. The log line now names the actual reason. Restoring the pooled stream and the prior profiling flag moves into an RAII guard that also captures the previous profiling state instead of hardcoding it back to enabled. Use MIOPEN_INTERNALS_EXPORT for StreamTracker, which requires rather than . Document the hipFreeAsync limitation: a skipped kernel keeps reading the caller's inputs, and the stream-ordered allocator does not wait on it. Co-Authored-By: Claude Opus 4.5 --- .../miopen/docs/reference/env_variables.rst | 8 ++ projects/miopen/src/conv/solver_finders.cpp | 102 +++++++++++++----- projects/miopen/src/hip/handlehip.cpp | 47 ++++++-- .../src/include/miopen/stream_tracker.hpp | 19 +++- projects/miopen/test/gtest/stream_tracker.cpp | 62 +++++++++++ 5 files changed, 198 insertions(+), 40 deletions(-) diff --git a/projects/miopen/docs/reference/env_variables.rst b/projects/miopen/docs/reference/env_variables.rst index 570c5655e567..2ddf801bec88 100644 --- a/projects/miopen/docs/reference/env_variables.rst +++ b/projects/miopen/docs/reference/env_variables.rst @@ -113,6 +113,14 @@ and :doc:`Performance database <../conceptual/perfdb>`. | 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. Freeing those tensors with + | ``hipFree`` is safe because it synchronizes the device first, + | but ``hipFreeAsync`` is ordered only against the stream it is + | given and can unmap memory the skipped kernel is still + | reading. Disable this option if the application frees + | convolution inputs through the stream-ordered allocator. - | 1 or unset: Enable (default) | 0: Disable diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index efa5b9c3cb0d..c9b723d8cb75 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -211,10 +211,68 @@ const std::vector>& GetConvSolverFinders() } // namespace conv -static float TryNaiveWithTimeout(const Handle& handle, - const Invoker& invoker, - const AnyInvokeParams& invoke_ctx, - float best_time) +namespace { + +struct NaiveWarmup +{ + enum class Status + { + Completed, + TimedOut, + ScratchUnavailable, + UnsupportedInvokeParams, + }; + + Status status; + float elapsed; +}; + +/// Redirects the handle onto a pooled stream with profiling off, restoring both +/// on scope exit so no early return can strand the handle off the root stream. +struct AutoPoolStream +{ + AutoPoolStream(const Handle& h, int pool_id) : handle(h), prev_profiling(h.IsProfilingEnabled()) + { + handle.SetStreamFromPool(pool_id); + handle.EnableProfiling(false); + } + + ~AutoPoolStream() + { + handle.SetStreamFromPool(0); + handle.EnableProfiling(prev_profiling); + } + + AutoPoolStream(const AutoPoolStream&) = delete; + AutoPoolStream& operator=(const AutoPoolStream&) = delete; + +private: + const Handle& handle; + bool prev_profiling; +}; + +std::string NaiveSkipReason(NaiveWarmup::Status status, float best_time) +{ + switch(status) + { + case NaiveWarmup::Status::TimedOut: + return "exceeded " + std::to_string(env::value(MIOPEN_NAIVE_TIMEOUT_FACTOR)) + + "% of best non-naive time (" + std::to_string(best_time) + " ms)"; + case NaiveWarmup::Status::ScratchUnavailable: + return "output tensor exceeds the scratch buffer cap"; + case NaiveWarmup::Status::UnsupportedInvokeParams: + return "invoke params type does not support scratch redirection"; + case NaiveWarmup::Status::Completed: break; + } + return "completed"; +} + +} // namespace + +static NaiveWarmup TryNaiveWithTimeout(const Handle& handle, + const Invoker& invoker, + const AnyInvokeParams& invoke_ctx, + float best_time) { std::shared_ptr scratch; AnyInvokeParams scratch_ctx; @@ -224,7 +282,7 @@ static float TryNaiveWithTimeout(const Handle& handle, auto params = invoke_ctx.CastTo(); scratch = handle.GetScratchBuffer(params.tensors.outDesc.GetNumBytes()); if(!scratch) - return -1.0f; + return {NaiveWarmup::Status::ScratchUnavailable, 0.0f}; params.tensors.out = scratch->buffer.get(); scratch_ctx = AnyInvokeParams{params}; } @@ -233,24 +291,18 @@ static float TryNaiveWithTimeout(const Handle& handle, auto params = invoke_ctx.CastTo(); scratch = handle.GetScratchBuffer(params.tensors.dwDesc.GetNumBytes()); if(!scratch) - return -1.0f; + return {NaiveWarmup::Status::ScratchUnavailable, 0.0f}; params.tensors.dw = scratch->buffer.get(); scratch_ctx = AnyInvokeParams{params}; } else - return -1.0f; + return {NaiveWarmup::Status::UnsupportedInvokeParams, 0.0f}; auto& tracker = handle.GetStreamTracker(); auto slot = tracker.acquire(handle); slot.scratch = scratch; - handle.SetStreamFromPool(slot.pool_id); - handle.EnableProfiling(false); - - auto guard = [&] { - handle.SetStreamFromPool(0); - handle.EnableProfiling(true); - }; + AutoPoolStream stream_guard{handle, slot.pool_id}; HipEventPtr ev_start = make_hip_event(); HipEventPtr ev_stop = make_hip_event(); @@ -267,7 +319,6 @@ static float TryNaiveWithTimeout(const Handle& handle, } catch(...) { - guard(); tracker.abandon(slot); throw; } @@ -288,18 +339,16 @@ static float TryNaiveWithTimeout(const Handle& handle, std::this_thread::yield(); } - guard(); - if(finished) { tracker.release(slot); float warmup_elapsed = 0.0f; (void)hipEventElapsedTime(&warmup_elapsed, ev_start.get(), ev_stop.get()); - return warmup_elapsed; + return {NaiveWarmup::Status::Completed, warmup_elapsed}; } tracker.abandon(slot); - return -1.0f; + return {NaiveWarmup::Status::TimedOut, 0.0f}; } /// Register invoker only for the best solution within algorithm. @@ -421,18 +470,17 @@ std::vector EvaluateInvokers(const Handle& handle, if(cutoff_naive) { - const float warmup = TryNaiveWithTimeout( + const auto warmup = TryNaiveWithTimeout( handle, invoker, invoke_ctx, core_result.find_search_best_time); - if(warmup < 0) + if(warmup.status != NaiveWarmup::Status::Completed) { - MIOPEN_LOG_I("Timeout: Naive Solver " - << algorithm_name.ToString() << ":" << sol.solver_id - << " exceeded " << env::value(MIOPEN_NAIVE_TIMEOUT_FACTOR) - << "% best budget (" << core_result.find_search_best_time - << " ms)"); + 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; + first_elapsed = warmup.elapsed; i = 1; } diff --git a/projects/miopen/src/hip/handlehip.cpp b/projects/miopen/src/hip/handlehip.cpp index 68045e3d14e4..f867f00045d7 100644 --- a/projects/miopen/src/hip/handlehip.cpp +++ b/projects/miopen/src/hip/handlehip.cpp @@ -426,22 +426,44 @@ std::shared_ptr Handle::GetScratchBuffer(std::size_t sz) cons return alloc; } -StreamTracker::Slot StreamTracker::acquire(const Handle& handle) +void StreamTracker::sweep() { - if(available_.empty()) + if(draining_.empty()) + return; + + for(auto it = draining_.begin(); it != draining_.end();) { - for(auto it = draining_.begin(); it != draining_.end();) + if(hipStreamQuery(it->stream) == hipSuccess) { - if(hipStreamQuery(it->stream) == hipSuccess) - { - it->scratch.reset(); - available_.push_back(std::move(*it)); - it = draining_.erase(it); - } - else - ++it; + 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()) { @@ -585,6 +607,9 @@ 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; diff --git a/projects/miopen/src/include/miopen/stream_tracker.hpp b/projects/miopen/src/include/miopen/stream_tracker.hpp index 56d38a64c22f..c28244582867 100644 --- a/projects/miopen/src/include/miopen/stream_tracker.hpp +++ b/projects/miopen/src/include/miopen/stream_tracker.hpp @@ -4,7 +4,7 @@ #ifndef GUARD_MIOPEN_STREAM_TRACKER_HPP_ #define GUARD_MIOPEN_STREAM_TRACKER_HPP_ -#include +#include #include #include @@ -21,7 +21,7 @@ struct ScratchAllocation std::size_t size = 0; }; -struct MIOPEN_EXPORT StreamTracker +struct MIOPEN_INTERNALS_EXPORT StreamTracker { struct Slot { @@ -30,8 +30,23 @@ struct MIOPEN_EXPORT StreamTracker 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(); diff --git a/projects/miopen/test/gtest/stream_tracker.cpp b/projects/miopen/test/gtest/stream_tracker.cpp index 5ec278505bfd..16fdb15ccf96 100644 --- a/projects/miopen/test/gtest/stream_tracker.cpp +++ b/projects/miopen/test/gtest/stream_tracker.cpp @@ -148,6 +148,68 @@ TEST_F(GPU_StreamTracker_FP32, CascadeAbandonReclaim) } } +TEST_F(GPU_StreamTracker_FP32, SweepReclaimsIdleStream) +{ + auto slot = tracker.acquire(handle); + int abandoned_id = slot.pool_id; + tracker.abandon(std::move(slot)); + + tracker.sweep(); + + auto reclaimed = tracker.acquire(handle); + EXPECT_EQ(reclaimed.pool_id, abandoned_id); + 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); + + int abandoned_id = slot.pool_id; + 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.pool_id, abandoned_id); + 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.pool_id, abandoned_id); + 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); From c41c531a3533ebf3ae794f3093ad4e4e72991ab1 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Wed, 12 Aug 2026 14:53:04 -0400 Subject: [PATCH 18/20] update wording for naive timout env limitation --- .../miopen/docs/reference/env_variables.rst | 17 +++++++++++------ 1 file changed, 11 insertions(+), 6 deletions(-) diff --git a/projects/miopen/docs/reference/env_variables.rst b/projects/miopen/docs/reference/env_variables.rst index 2ddf801bec88..c505e2d97967 100644 --- a/projects/miopen/docs/reference/env_variables.rst +++ b/projects/miopen/docs/reference/env_variables.rst @@ -115,12 +115,17 @@ and :doc:`Performance database <../conceptual/perfdb>`. | succeeds. | | Limitation: a skipped kernel keeps reading the caller's input - | tensors until it finishes. Freeing those tensors with - | ``hipFree`` is safe because it synchronizes the device first, - | but ``hipFreeAsync`` is ordered only against the stream it is - | given and can unmap memory the skipped kernel is still - | reading. Disable this option if the application frees - | convolution inputs through the stream-ordered allocator. + | 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 From cd4acb0ac760da9e92b45f7d0da1b5916acc70f6 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Wed, 12 Aug 2026 21:35:33 -0400 Subject: [PATCH 19/20] fix(miopen): use tracker-owned streams for naive timeout evaluation StreamTracker::acquire() grew the Handle's index-addressed stream pool and handed back the pool id. MHA and RNN claim ids in that same pool by hardcoded number, so a tracker slot could alias a stream another solver was already writing to. Reserving an id also eagerly constructed every stream below it along with its rocBLAS and hipBLASLt handles. The tracker now creates and owns its streams via Handle::CreateExclusiveStream and addresses slots by stream pointer rather than pool id. Handle gains a thread-local exclusive-stream override that GetStream() honors ahead of the pool index; rhandle() and HipblasLtHandle() throw while it is set, since those handles stay bound to the pool stream and would otherwise submit work to the wrong queue. Owned streams are declared before the slot vectors so they outlive the slots that borrow them and are destroyed after the destructor synchronizes. Also corrects the meopenHandle_current_stream_id spelling. --- projects/miopen/src/conv/solver_finders.cpp | 21 +++--- projects/miopen/src/hip/handlehip.cpp | 50 +++++++++---- projects/miopen/src/include/miopen/handle.hpp | 11 +++ .../src/include/miopen/stream_tracker.hpp | 14 +++- projects/miopen/test/gtest/stream_tracker.cpp | 71 ++++++++++++------- 5 files changed, 112 insertions(+), 55 deletions(-) diff --git a/projects/miopen/src/conv/solver_finders.cpp b/projects/miopen/src/conv/solver_finders.cpp index c9b723d8cb75..c3b3d160a0c0 100644 --- a/projects/miopen/src/conv/solver_finders.cpp +++ b/projects/miopen/src/conv/solver_finders.cpp @@ -227,24 +227,25 @@ struct NaiveWarmup float elapsed; }; -/// Redirects the handle onto a pooled stream with profiling off, restoring both -/// on scope exit so no early return can strand the handle off the root stream. -struct AutoPoolStream +/// 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 { - AutoPoolStream(const Handle& h, int pool_id) : handle(h), prev_profiling(h.IsProfilingEnabled()) + AutoExclusiveStream(const Handle& h, hipStream_t stream) + : handle(h), prev_profiling(h.IsProfilingEnabled()) { - handle.SetStreamFromPool(pool_id); + handle.SetExclusiveStream(stream); handle.EnableProfiling(false); } - ~AutoPoolStream() + ~AutoExclusiveStream() { - handle.SetStreamFromPool(0); + handle.SetExclusiveStream(nullptr); handle.EnableProfiling(prev_profiling); } - AutoPoolStream(const AutoPoolStream&) = delete; - AutoPoolStream& operator=(const AutoPoolStream&) = delete; + AutoExclusiveStream(const AutoExclusiveStream&) = delete; + AutoExclusiveStream& operator=(const AutoExclusiveStream&) = delete; private: const Handle& handle; @@ -302,7 +303,7 @@ static NaiveWarmup TryNaiveWithTimeout(const Handle& handle, auto slot = tracker.acquire(handle); slot.scratch = scratch; - AutoPoolStream stream_guard{handle, slot.pool_id}; + AutoExclusiveStream stream_guard{handle, slot.stream}; HipEventPtr ev_start = make_hip_event(); HipEventPtr ev_stop = make_hip_event(); diff --git a/projects/miopen/src/hip/handlehip.cpp b/projects/miopen/src/hip/handlehip.cpp index f867f00045d7..dd771340ee61 100644 --- a/projects/miopen/src/hip/handlehip.cpp +++ b/projects/miopen/src/hip/handlehip.cpp @@ -171,7 +171,10 @@ int set_default_device() #endif // NOLINTNEXTLINE (cppcoreguidelines-avoid-non-const-global-variables) -static thread_local unsigned int meopenHandle_current_stream_id = 0; +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; @@ -290,7 +293,7 @@ struct HandleImpl Handle::Handle(miopenAcceleratorQueue_t stream) : impl(std::make_unique()) { - meopenHandle_current_stream_id = 0; + miopenHandle_current_stream_id = 0; this->impl->device = get_device_id(); if(stream == nullptr) @@ -317,7 +320,7 @@ Handle::Handle(miopenAcceleratorQueue_t stream) : impl(std::make_unique()) { - meopenHandle_current_stream_id = 0; + miopenHandle_current_stream_id = 0; #if MIOPEN_BUILD_DEV this->impl->device = set_default_device(); this->impl->root_stream = impl->create_stream(); @@ -348,7 +351,7 @@ Handle::~Handle() {} // not MT safe void Handle::SetStream(miopenAcceleratorQueue_t streamID) const { - meopenHandle_current_stream_id = 0; + miopenHandle_current_stream_id = 0; this->impl->root_stream = HandleImpl::reference_stream(streamID); @@ -363,7 +366,7 @@ void Handle::SetStream(miopenAcceleratorQueue_t streamID) const MIOPEN_LOG_NQI(*this); } -void Handle::SetStreamFromPool(int streamID) const { meopenHandle_current_stream_id = streamID; } +void Handle::SetStreamFromPool(int streamID) const { miopenHandle_current_stream_id = streamID; } void Handle::ReserveExtraStreamsInPool(int cnt) const { @@ -394,13 +397,26 @@ 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(meopenHandle_current_stream_id == 0) + 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 std::shared_lock lock(this->impl->stream_pool_mutex); - return this->impl->ms_resourse_ptr->stream_pool.at(meopenHandle_current_stream_id - 1).get(); + return this->impl->ms_resourse_ptr->stream_pool.at(miopenHandle_current_stream_id - 1).get(); } StreamTracker& Handle::GetStreamTracker() const { return impl->stream_tracker_; } @@ -472,10 +488,8 @@ StreamTracker::Slot StreamTracker::acquire(const Handle& handle) return slot; } - int id = next_id_++; - handle.ReserveExtraStreamsInPool(id); - handle.SetStreamFromPool(id); - return {id, handle.GetStream(), {}}; + owned_streams_.push_back(handle.CreateExclusiveStream()); + return {owned_streams_.back().get(), {}}; } void Handle::SetAllocator(miopenAllocatorFunction allocator, @@ -996,11 +1010,14 @@ Handle::CreateSubBuffer(ConstData_t data, std::size_t offset, std::size_t size) const rocblas_handle_ptr& Handle::rhandle() const { - if(meopenHandle_current_stream_id == 0) + 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 std::shared_lock lock(this->impl->stream_pool_mutex); - return this->impl->ms_resourse_ptr->rhandle_pool.at(meopenHandle_current_stream_id - 1); + return this->impl->ms_resourse_ptr->rhandle_pool.at(miopenHandle_current_stream_id - 1); } rocblas_handle_ptr Handle::CreateRocblasHandle(miopenAcceleratorQueue_t stream) const @@ -1016,11 +1033,14 @@ rocblas_handle_ptr Handle::CreateRocblasHandle(miopenAcceleratorQueue_t stream) #if MIOPEN_USE_HIPBLASLT const hipblasLt_handle_ptr& Handle::HipblasLtHandle() const { - if(meopenHandle_current_stream_id == 0) + 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 std::shared_lock lock(this->impl->stream_pool_mutex); - return this->impl->ms_resourse_ptr->hhandle_pool.at(meopenHandle_current_stream_id - 1); + return this->impl->ms_resourse_ptr->hhandle_pool.at(miopenHandle_current_stream_id - 1); } hipblasLt_handle_ptr Handle::CreateHipblasLtHandle() const diff --git a/projects/miopen/src/include/miopen/handle.hpp b/projects/miopen/src/include/miopen/handle.hpp index 836454c97ebf..71ae1f5e34fd 100644 --- a/projects/miopen/src/include/miopen/handle.hpp +++ b/projects/miopen/src/include/miopen/handle.hpp @@ -95,6 +95,17 @@ 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, diff --git a/projects/miopen/src/include/miopen/stream_tracker.hpp b/projects/miopen/src/include/miopen/stream_tracker.hpp index c28244582867..3d36897a589e 100644 --- a/projects/miopen/src/include/miopen/stream_tracker.hpp +++ b/projects/miopen/src/include/miopen/stream_tracker.hpp @@ -8,6 +8,7 @@ #include #include +#include #include #include @@ -21,12 +22,17 @@ struct ScratchAllocation 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 { - int pool_id; - hipStream_t stream; + hipStream_t stream = nullptr; std::shared_ptr scratch; }; @@ -56,9 +62,11 @@ struct MIOPEN_INTERNALS_EXPORT StreamTracker 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_; - int next_id_ = 1; }; } // namespace miopen diff --git a/projects/miopen/test/gtest/stream_tracker.cpp b/projects/miopen/test/gtest/stream_tracker.cpp index 16fdb15ccf96..ba0145c71ba6 100644 --- a/projects/miopen/test/gtest/stream_tracker.cpp +++ b/projects/miopen/test/gtest/stream_tracker.cpp @@ -11,6 +11,7 @@ #include #include +#include namespace { @@ -49,22 +50,41 @@ class GPU_StreamTracker_FP32 : public ::testing::Test TEST_F(GPU_StreamTracker_FP32, AcquireRelease) { auto slot = tracker.acquire(handle); - ASSERT_GT(slot.pool_id, 0); ASSERT_NE(slot.stream, nullptr); - int saved_id = slot.pool_id; + auto saved_stream = slot.stream; tracker.release(slot); auto slot2 = tracker.acquire(handle); - EXPECT_EQ(slot2.pool_id, saved_id); + 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.pool_id, slot2.pool_id); + EXPECT_NE(slot1.stream, slot2.stream); tracker.release(slot2); tracker.release(slot1); @@ -78,13 +98,13 @@ TEST_F(GPU_StreamTracker_FP32, AbandonAndReclaim) ASSERT_EQ(hipMalloc(&dev_ptr, 64), hipSuccess); ASSERT_EQ(hipMemsetAsync(dev_ptr, 0, 64, slot.stream), hipSuccess); - int abandoned_id = slot.pool_id; + auto abandoned_stream = slot.stream; tracker.abandon(slot); - ASSERT_EQ(hipStreamSynchronize(slot.stream), hipSuccess); + ASSERT_EQ(hipStreamSynchronize(abandoned_stream), hipSuccess); auto reclaimed = tracker.acquire(handle); - EXPECT_EQ(reclaimed.pool_id, abandoned_id); + EXPECT_EQ(reclaimed.stream, abandoned_stream); tracker.release(reclaimed); ASSERT_EQ(hipFree(dev_ptr), hipSuccess); @@ -97,18 +117,18 @@ TEST_F(GPU_StreamTracker_FP32, AbandonStillDraining) StreamGate gate; ASSERT_EQ(hipLaunchHostFunc(slot.stream, StreamGate::callback, &gate), hipSuccess); - int abandoned_id = slot.pool_id; + auto abandoned_stream = slot.stream; tracker.abandon(slot); auto next = tracker.acquire(handle); - EXPECT_NE(next.pool_id, abandoned_id); + EXPECT_NE(next.stream, abandoned_stream); gate.open(); - ASSERT_EQ(hipStreamSynchronize(slot.stream), hipSuccess); + 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.pool_id, abandoned_id); + EXPECT_EQ(reclaimed.stream, abandoned_stream); tracker.release(reclaimed); tracker.release(next); } @@ -117,47 +137,45 @@ TEST_F(GPU_StreamTracker_FP32, CascadeAbandonReclaim) { constexpr int kCount = 4; std::vector slots; - std::vector abandoned_ids; + std::set seen; for(int i = 0; i < kCount; ++i) { auto slot = tracker.acquire(handle); - abandoned_ids.push_back(slot.pool_id); + seen.insert(slot.stream); tracker.abandon(slot); } for(int i = 0; i < kCount; ++i) { - slots.emplace_back(tracker.acquire(handle)); + auto slot = tracker.acquire(handle); + seen.insert(slot.stream); + slots.emplace_back(std::move(slot)); } - // Wait for everything to drain for(auto& s : slots) tracker.release(s); - // Now acquire kCount — all should come from available (no new pool growth) - auto before = tracker.acquire(handle); - int max_id = before.pool_id; - tracker.release(before); - + // 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_LE(s.pool_id, max_id); + EXPECT_EQ(seen.count(s.stream), 1u); tracker.release(s); } } TEST_F(GPU_StreamTracker_FP32, SweepReclaimsIdleStream) { - auto slot = tracker.acquire(handle); - int abandoned_id = slot.pool_id; + 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.pool_id, abandoned_id); + EXPECT_EQ(reclaimed.stream, abandoned_stream); tracker.release(reclaimed); } @@ -168,7 +186,6 @@ TEST_F(GPU_StreamTracker_FP32, SweepLeavesBusyStreamDraining) StreamGate gate; ASSERT_EQ(hipLaunchHostFunc(slot.stream, StreamGate::callback, &gate), hipSuccess); - int abandoned_id = slot.pool_id; auto busy_stream = slot.stream; tracker.abandon(std::move(slot)); @@ -176,7 +193,7 @@ TEST_F(GPU_StreamTracker_FP32, SweepLeavesBusyStreamDraining) // Still gated, so the slot must not have been reclaimed auto next = tracker.acquire(handle); - EXPECT_NE(next.pool_id, abandoned_id); + EXPECT_NE(next.stream, busy_stream); tracker.release(next); gate.open(); @@ -186,7 +203,7 @@ TEST_F(GPU_StreamTracker_FP32, SweepLeavesBusyStreamDraining) // available_ is LIFO, so the just-swept slot is on top auto reclaimed = tracker.acquire(handle); - EXPECT_EQ(reclaimed.pool_id, abandoned_id); + EXPECT_EQ(reclaimed.stream, busy_stream); tracker.release(reclaimed); } From e1a5e5c91e44359551b102dec0f8a6b2e1532a21 Mon Sep 17 00:00:00 2001 From: Christopher Erb Date: Fri, 21 Aug 2026 15:38:01 -0400 Subject: [PATCH 20/20] test(miopen): run naive-timeout gtests in CI and skip winograd-only archs The StreamTracker and NaiveTimeout suites are plain TEST/TEST_F with no INSTANTIATE_TEST_SUITE_P, so they carry no "Prefix/" and none of the existing "*/GPU_*" patterns in test_categories.yaml could match them. 19 of the 20 new tests were silently never executed; only the CPU case ran, via the blanket "CPU_*" pattern. Add unprefixed patterns so the standard/comprehensive/full suites pick all 20 up. Now that they actually run, guard the two assertions that are only meaningful where ConvDirectNaive is a candidate. On gfx803/gfx900/ gfx906/gfx908, ConvBinWinograd3x3U is applicable to the test's convolution, so IsWinograd3x3SupportedAndFast() sets use_winograd_only and the Direct finder is disabled outright -- naive presence or absence then says nothing about the timeout. Co-Authored-By: Claude Sonnet 4 --- projects/miopen/test/gtest/naive_timeout.cpp | 19 +++++++++++++++++++ .../miopen/test/gtest/test_categories.yaml | 4 ++++ 2 files changed, 23 insertions(+) diff --git a/projects/miopen/test/gtest/naive_timeout.cpp b/projects/miopen/test/gtest/naive_timeout.cpp index a61670e8772a..302462c4208e 100644 --- a/projects/miopen/test/gtest/naive_timeout.cpp +++ b/projects/miopen/test/gtest/naive_timeout.cpp @@ -98,6 +98,17 @@ std::vector RunFind(miopenHandle_t handle, size_t max_solutions = 8 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 four +// 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 == "gfx803" || name == "gfx900" || name == "gfx906" || name == "gfx908"; +} + } // namespace // With MIOPEN_NAIVE_TIMEOUT=0 (opt-out), a naive solver must appear somewhere in the @@ -107,6 +118,10 @@ 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); @@ -165,6 +180,10 @@ 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); diff --git a/projects/miopen/test/gtest/test_categories.yaml b/projects/miopen/test/gtest/test_categories.yaml index 268ddd1579f9..ab22966c6f22 100644 --- a/projects/miopen/test/gtest/test_categories.yaml +++ b/projects/miopen/test/gtest/test_categories.yaml @@ -64,6 +64,10 @@ 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