Skip to content

[MIOPEN] gfx950 and gfx942 heuristics rules for solver selection - #7780

Closed
jdcampbe wants to merge 6 commits into
developfrom
users/jascampb/miopen-heuristics-rules
Closed

jdcampbe wants to merge 6 commits into
developfrom
users/jascampb/miopen-heuristics-rules

Conversation

@jdcampbe

Copy link
Copy Markdown
Contributor

Motivation

MIOPEN_FIND_MODE=2 (Fast/Immediate mode) currently dispatches all
solver selection through the TunaNet AI model in
ai::immed_mode::PredictSolver. Offline analysis of the gfx950 and
gfx942 perf databases shows there are large regions of problem space
where a small, deterministic rules-based dispatcher matches or beats
TunaNet's pick — and where TunaNet's mispredictions cost meaningful
runtime. This PR wires up per-architecture rules-based pickers that
short-circuit TunaNet when they have a confident answer, and fall
through to TunaNet otherwise. Behavior on non-gfx942/gfx950 devices is
unchanged.

Technical Details

Adds two new translation units, modifies the immediate-mode entry
point, and registers them in CMake. All code is guarded by the
existing MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK build flag.

New files

  • projects/miopen/src/include/miopen/conv/heuristics/gfx950_rules.hpp
  • gfx942_rules.hpp — single-function headers declaring solver::Id
    PickSolver(const conv::ProblemDescription&) in miopen::ai::gfx950 /
    miopen::ai::gfx942 namespaces.
  • projects/miopen/src/conv/heuristics/gfx950_rules.cpp (~280 lines)
    and gfx942_rules.cpp (~944 lines) — per-architecture rules
    dispatchers. Each is a pure function: a cached SolverIds struct
    constructed once on first call, then a tree of per-direction /
    per-dtype helpers (PickFwdFp32/Fp16/Bfp16, PickBwdData, PickWrW,
    PickSolver3d) dispatched from a single top-level PickSolver. 3D and
    transposed-convolution problems return an invalid Id from gfx950
    (deferred to TunaNet).

Modified

  • projects/miopen/src/conv/heuristics/ai_heuristics.cpp — inside
    PredictSolver, immediately after the prediction-cache check, two
    parallel short-circuit blocks (gfx950, gfx942). Each:
    a. Skips if the device doesn't match or the corresponding
    MIOPEN_DEBUG_GFX{950,942}_RULES_PICK=0 opt-out is set.
    b. Calls the rules picker. If it returns a valid solver::Id and that
    solver's IsApplicable(ctx, problem) is true, stores the pick in
    PredictionCache and returns. Otherwise logs the reason and falls
    through to the existing TunaNet path.
    c. Logging discriminates three outcomes:
    - "rules picked " — successful short-circuit.
    • "rules: no pick (problem outside rules coverage), falling
      through to TunaNet" — picker returned {} (3D / transposed / unhandled
      dtype / no rule matched).
    • "rules picked but solver reported not applicable,
      falling through to TunaNet" — picker chose a solver, but IsApplicable
      rejected it (workspace/dim/dtype gate).
  • projects/miopen/src/CMakeLists.txt — registers the two new .cpp
    files inside the existing MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK block.

Test Plan

  1. In-container per-TU compile. Built a miopen-deps:local image from
    projects/miopen/Dockerfile (target miopen_deps) and ran the exact
    compile command from compile_commands.json for each of the three
    changed TUs (gfx950_rules.cpp, gfx942_rules.cpp, ai_heuristics.cpp).
    Backend: HIPNOGPU, MIOPEN_USE_COMPOSABLEKERNEL=Off,
    MIOPEN_USE_AI_KERNEL_TUNING_NET=Off.
  2. End-to-end smoke on gfx950. Ran a real-workload conv driver
    coverage trace with MIOPEN_FIND_MODE=2 MIOPEN_LOG_LEVEL=4 on gfx950
    and grepped the resulting log for the new log lines.

Test Result

  • Per-TU compile: all three changed TUs exit 0.
  • End-to-end gfx950 trace
    • 86 gfx950 rules lines total.
    • 32 successful picks (all resolved to
      ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC in this workload).
    • 54 fall-throughs to TunaNet (the new discriminated logging will
      tell us in the next run how many of those are out-of-coverage vs.
      not-applicable).

Risks

  • Wrong pick can hurt latency. If the rules pick a solver that is
    technically applicable but slower than what TunaNet would have picked,
    that's a silent perf regression on Fast mode for affected shapes.
    Mitigation: MIOPEN_DEBUG_GFX950_RULES_PICK=0 /
    MIOPEN_DEBUG_GFX942_RULES_PICK=0 instantly reverts each arch to
    TunaNet-only behavior, no rebuild required.
  • Behavior change is scoped, but Fast-mode users see it immediately.
    The change has zero effect on MIOPEN_FIND_MODE ≠ 2 and zero effect on
    non-gfx942/gfx950 devices. But on gfx942/gfx950 in Fast mode, this is
    on by default — early adopters of those archs in Fast mode will get
    the new dispatcher without opting in. The env-var opt-out is the
    escape hatch.
  • Rules may diverge from reality as the perf-DB and solver set evolve.
    Rules were derived from a perf-DB snapshot; new solvers, removed
    solvers, or changed solver applicability can leave the rules picking
    stale targets. The not-applicable log path catches the worst case
    (picked-but-rejected) silently; the bigger risk is picking an
    applicable-but-suboptimal solver, which is invisible from logs.
    Refresh cadence and a telemetry hook capturing the actual post-search
    winner are tracked separately (out of scope for this PR).
  • Solver-ID lookups happen at first call. Each SolverIds struct does
    ~15–25 solver::Id{name} lookups on the first invocation. A typo in a
    solver name would surface as an invalid Id and silently degrade the
    rule to "no pick."

Submission Checklist

jdcampbe and others added 3 commits May 24, 2026 09:25
When MIOPEN_FIND_MODE=2 on gfx950, short-circuit the Tunanet ND model
with a deterministic rules-based dispatcher derived from gfx950 perf-DB
analysis (val geomean_slowdown = 1.0171, 0 invalid picks on 30k val
problems). Falls through to Tunanet on miss (3D, transposed conv, int8,
or no applicable solver). Gated by MIOPEN_DEBUG_GFX950_RULES_PICK
(on by default).

Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
Mirrors the existing gfx950 rules hook. When MIOPEN_FIND_MODE=2 on
gfx942, short-circuit the Tunanet ND model with a deterministic
dispatcher ported from /home/jascampb/AutoResearchGfx942/rules.py.
Falls through to Tunanet when no rule applies. Covers 2D fwd/bwd/wrw
across fp32/fp16/bf16/int8, 3D dispatch, and transposed conv. Gated by
MIOPEN_DEBUG_GFX942_RULES_PICK (on by default).

Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
…n logs

Split the single "no applicable pick" log line into two cases for both
gfx950 and gfx942 rules hooks: rules returned an invalid Id (problem
outside coverage) vs. rules picked a solver that reported not
applicable. Lets us tell from logs why fall-through to TunaNet happened.

Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
Rewrites the gfx950 rules dispatcher as a flat carve-out sequence
mirroring the regenerated rules from the external mining project,
adding 3D group-conv routing and iter-tagged shape carve-outs.
Comment on lines +751 to +800
// gfx950 rules-based short-circuit: deterministic dispatcher derived from
// perf-DB analysis; preempts the ND TunaNet model when applicable.
if(device == "gfx950" && !env::disabled(MIOPEN_DEBUG_GFX950_RULES_PICK))
{
const auto picked = gfx950::PickSolver(problem);
if(!picked.IsValid())
{
MIOPEN_LOG_I2("gfx950 rules: no pick (problem outside rules coverage), "
"falling through to TunaNet");
}
else if(!picked.GetSolver().IsApplicable(ctx, problem))
{
MIOPEN_LOG_I2("gfx950 rules picked " << picked.ToString()
<< " but solver reported not applicable, "
"falling through to TunaNet");
}
else
{
MIOPEN_LOG_I2("gfx950 rules picked " << picked.ToString());
std::vector<std::any> any_sol;
any_sol.push_back(picked.Value());
StorePredictionCache(problem, device, any_sol);
return {picked.Value()};
}
}

// gfx942 rules-based short-circuit: same pattern as gfx950 above.
if(device == "gfx942" && !env::disabled(MIOPEN_DEBUG_GFX942_RULES_PICK))
{
const auto picked = gfx942::PickSolver(problem);
if(!picked.IsValid())
{
MIOPEN_LOG_I2("gfx942 rules: no pick (problem outside rules coverage), "
"falling through to TunaNet");
}
else if(!picked.GetSolver().IsApplicable(ctx, problem))
{
MIOPEN_LOG_I2("gfx942 rules picked " << picked.ToString()
<< " but solver reported not applicable, "
"falling through to TunaNet");
}
else
{
MIOPEN_LOG_I2("gfx942 rules picked " << picked.ToString());
std::vector<std::any> any_sol;
any_sol.push_back(picked.Value());
StorePredictionCache(problem, device, any_sol);
return {picked.Value()};
}
}

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

These are identical, and I think we could mostly merge this into a single re-usable method that uses the device details for deciding which rules to use.

*******************************************************************************/

// Source of truth for the decision tree:
// /home/jascampb/AutoResearchGfx942/rules.py

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We should include this script, or remove comment.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this an auto-generated decision tree?

Looks like a codified criteria matcher, and could be potentially generalized, but might not be worth the effort.

jdcampbe and others added 2 commits May 28, 2026 12:48
Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
…y + abstain)

Full re-port against the 4146-line rules.py refresh:
- Base layer (PickBase): drops deprecated iter167/163/170 NAIVE_FWD routes
  (now regress post-data-refresh), adds 2D carve-backs iter258/260/261/262/
  263/259, the iter266 OR-clause, the 3x3 s=1 squeeze→Wino pair iter264/265,
  the odd-channel WRW iter257, and the real-3D fz>=2 / depth>=2 / first-conv-
  3D-stride-2 rule suite (iter268/269/273/274/277).
- Overlay layer (ApplyOverlay): post-base rewrites — iter276/272/268/268b/
  269b/270/271/271b/278-282/289-303/304-332. First match wins.
- Abstain layer (ShouldAbstain): whole-cohort (iter337-339) and NCHW sub-
  cohort (iter340-345) entries that defer to TunaNet by returning {}.
  NaN-layout entries are omitted (MIOpen always populates in_layout).

PickSolver is now a thin orchestrator: guards → PickBase → ApplyOverlay →
abstain check.

Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants