Repository navigation
feat(miopen): Credit TransposingSolver wrapper in performance logs - #10649
Merged
Merged
Conversation
Record solver name and invoker timing for ConvolutionForwardBias and ConvolutionBackwardBias paths in conv_driver so bias kernels appear in performance logs alongside the main convolution kernels. Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
When MIOPEN_PERFORMANCE_LOGS>0 is set, the JSON records emitted via
LogSolutionName are consumed by the perf-eval pipeline and parsed into
the conv_run / conv_problem / performance_config tables that downstream
SystemDB generation and ML decision-tree training read from.
A TransposingSolver wrapper (e.g. TransposedConvBinWinoRxS<2-3> wrapping
the NHWC-native ConvBinWinogradRxSf2x3) overrides IsApplicable to return
true on NCHW inputs, then delegates Search/GetSolution to the inner
solver against a transposed NHWC view of the problem. The inner's own
GenericSearch loop calls LogSolutionName with the inner's SolverDbId(),
so the perf log credits the inner against the user's original NCHW
problem key -- even though the actual runtime candidate is the wrapper
and the inner's IsApplicable would correctly reject a uniform-NCHW
problem if asked directly.
Downstream this surfaces in two ways:
- The generated FDB carries the inner solver's name against the NCHW
find_dbkey. dbsync's StaticFDBSync test calls IsApplicable on that
inner and gets `false`, flagging the entry as NotApplicable.
- The tuned PerformanceConfig is keyed under the NHWC perf_dbkey, so
a runtime lookup keyed by the NCHW problem cannot find it.
Comparing the generated FDB against a gathered (MIOpen-direct) UDB on
gfx950 showed 61% winner disagreement, 68% of which were the
NHWC-inner-on-NCHW-key pattern. Same pattern on gfx942 at 38%/32%. ML
convergence on gfx942 specifically is poor and the mislabeling is the
strongest hypothesis for the gap with gfx950.
Fix: add SolverNameOverride (thread-local) plus a ScopedSolverNameOverride
RAII helper to kernel_tuning_mode.hpp, and apply it in
TransposingSolver::Search and both GetSolution specializations. When the
inner's logging path subsequently calls LogSolutionName(inner_name, ...),
the override substitutes the wrapper's name + id, so the JSON record is
attributed to the registered candidate the runtime would actually invoke.
Zero overhead when MIOPEN_PERFORMANCE_LOGS=0 -- the override is set
inside the scope but LogSolutionName short-circuits at IsLoggingKernel()
before reading it.
Wrap the SolverDbId()/Id::Value() lookups and the thread-local override
state-touches behind IsPerformanceLoggingEnabled() in both
TransposingSolver::Search and both GetSolution specializations, plus
make ScopedSolverNameOverride's ctor/dtor early-exit when logging is off.
Net effect when MIOPEN_PERFORMANCE_LOGS=0 (normal production):
- one env::value() read per Search/GetSolution call (already cached
after the first read);
- one branch;
- no SolverDbId() call, no solver::Id lookup, no thread-local write.
When MIOPEN_PERFORMANCE_LOGS>0 (perf-eval only): identical behavior to
the previous commit -- the override is set, LogSolutionName honors it,
and the JSON record is attributed to the wrapper instead of the inner.
jdcampbe
changed the title
jdcampbe
marked this pull request as ready for review
August 12, 2026 00:19
Contributor
Author
Compile check: PASSHIPNOGPU Docker build (
Both changed headers confirmed present in Not verified
|
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Collapse three-line ternary initializers that fit on one line, as clang-format 18.1.4 (the pre-commit pinned version) formats them. Whitespace only; no behavior change. Co-Authored-By: Claude <noreply@anthropic.com>
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Sep 2, 2026
## Summary Backfills the **MIOpen 3.6.1 / ROCm 10.1.0** changelog section with externally-visible changes that landed since the `release/therock-10.0` branch but weren't yet documented. All 43 MIOpen commits in that range were reviewed; CI/build-infra, test-only, docs, version-bump, and already-documented changes were excluded. 21 entries added: ### Added * [RNN] Fused single-layer LSTM forward-inference kernel, gated behind `MIOPEN_DEBUG_RNN_FUSED_INFERENCE` (#8915) * [Conv] Naive reference kernel enabling 3D INT8 forward convolution (#10051) * [Conv] `ConvHipConv` solver backed by vendored hipconv kernels (#9746) * [Conv] float8 vector type support + deduplicated math kernels (#10130) * [Conv] Validated gfx1151 and gfx1200 (Navi) SystemDBs (#10050) * [Conv] `MIOPEN_DEBUG_DISABLE_SYSTEM_DB` / `MIOPEN_DEBUG_DISABLE_USER_DB` env vars (#10728) * [Conv] CK depthwise convolution on RDNA wave32 GPUs (#9373) ### Changed * [Conv] Winograd Rage RxS f2x3 on gfx950 + gfx942 v4_6_1/v4_9_1 kernels (#9014) * [Conv] Updated 3D AI heuristics for gfx942 & gfx950 (#8954) ### Optimized * [Conv] Single GEMM for point (1x1) output convolution, all directions, 2D/3D (#10912) * [Conv] Naive backward-data kernel work group widened to 1024 (#10684) ### Resolved Issues * [Conv] Batch-split miscalculation on large WrW tensors — ROCM-26011 (#11188) * [Conv] Conv failure on XNACK-unsupported gfx9 APUs (gfx90c) (#11211) * [BatchNorm] Device heap-overflow in fwd spatial `FinalMeanVariance` — ALMIOPEN-2424 (#11093) * [Conv] Deterministic `SetNextValue` polarity for grouped Bwd/Wrw CK solvers (#10240) * [LRN][LayerNorm] Grid alignment + OOB threads on gfx1250 (#10251) * [Conv] Stale gfx950 grouped-wrw SystemDB entries (#11130) * [Conv][Windows] System-DB path resolution from the MIOpen module (#10838) * [Conv] System-DB fallback so MI308X (gfx942) resolves a usable DB (#8802) * JSON perf logs (`MIOPEN_PERFORMANCE_LOGS`) dropping the last solver during Find (#11166) * Build failure against GCC 15 libstdc++ under C++20 (#9642) ## Notes Deliberately excluded: internal/default-off changes (#10219, #10088, #10649, #10231), a near-duplicate of an existing 10.0.0 entry (#10188), and CI/build-infra/test-only/docs commits. Docs-only change; no code impact. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : ALMIOPEN-2426
Motivation
When
MIOPEN_PERFORMANCE_LOGS>0, the JSON records emitted viaLogSolutionNameareconsumed by the perf-eval pipeline and parsed into the
conv_run/conv_problem/performance_configtables that downstream SystemDB generation and ML decision-treetraining read from.
A
TransposingSolverwrapper delegatesSearch/GetSolutionto an inner solver againsta transposed view of the problem. The inner's own
GenericSearchloop callsLogSolutionNamewith the inner'sSolverDbId(), so the perf log credits the inneragainst the user's original problem key — even though the runtime candidate is the
wrapper, and the inner would correctly reject that problem if asked directly.
Downstream this surfaces two ways:
dbsync's
StaticFDBSynccallsIsApplicableon that inner, getsfalse, and flagsthe entry
NotApplicable.PerformanceConfigis keyed under the inner's perf_dbkey, so a runtimelookup keyed by the original problem cannot find it.
Technical Details
kernel_tuning_mode.hpp: add a thread-localSolverNameOverrideplus aScopedSolverNameOverrideRAII helper.transposing_solver.hpp: apply it inTransposingSolver::Searchand bothGetSolutionspecializations (3 call sites), so when the inner's logging path callsLogSolutionName(inner_name, ...)the override substitutes the wrapper's name and id.SolverDbId()/solver::Idlookups and the thread-local writes are gated behindIsPerformanceLoggingEnabled(). WithMIOPEN_PERFORMANCE_LOGS=0(normal production)the cost is one cached env-var read and one branch — no
SolverDbId()call, noregistry lookup, no thread-local write.
conv_driver.hpp: performance logging for bias convolution kernels.Scope — what this does not change
This is perf-log attribution only, and it is inert unless
MIOPEN_PERFORMANCE_LOGS>0.It does not change solver selection, applicability, or any runtime find behavior.
In particular it does not address the separate issue where
Transposed*solvers areskipped during
DYNAMIC_HYBRIDfind becauseTransposingSolver::IsDynamic()returns ahardcoded
false— that is a distinct defect on the runtime discovery path and needs itsown fix.
Risk Assessment
Risk Level: Low. Behavior is unchanged when performance logging is off, which is every
production path. The affected code is confined to the transposing-solver wrapper and a new
RAII helper.
Test Plan
and confirm
Transposed*wrappers are credited in place of their inner solvers, andthat the FDB/UDB winner-disagreement rate drops.
Co-authored-by: Claude noreply@anthropic.com
🤖 Generated with Claude Code