Skip to content

fix(miopen): read ConvHipConv perf config descriptor whole; drop it from the pcfg bundle - #13335

Closed
jdcampbe wants to merge 2 commits into
developfrom
users/jascampb/hipconv-pcfg-serialization-fix
Closed

jdcampbe wants to merge 2 commits into
developfrom
users/jascampb/hipconv-pcfg-serialization-fix

Conversation

@jdcampbe

@jdcampbe jdcampbe commented Oct 8, 2026 •

Copy link
Copy Markdown
Contributor

Summary

Fixes a serialization bug in ConvHipConv's performance config and removes the solver from the shipped perf-config bundle until its model is retrained.

PerformanceConfigConvHipConv stores a kernel descriptor string, e.g. direct[tile_size_k=128,tile_size_n=1,tile_size_h=16,tile_size_w=16,aligned=false]. #9680 switched the config to this string form but left it on PerfConfigBase, whose field reader splits on commas. So on read the descriptor was truncated to direct[tile_size_k=128. Consequences:

  • the descriptor matched no kernel, so the perf-config picker skipped every ConvHipConv candidate; and
  • a tuned ConvHipConv config written to the perf-db could not be reloaded (it deserialized to the truncated, unresolvable value).

The comma format is not a hipconv defect — describe_config() has always produced family[field=value,...] (used by perf logs, #11636). #9680 adopted it as the saved form without making the reader handle commas.

JIRA ID : ALMIOPEN-2494

Changes

  • fix: read ConvHipConv perf config descriptor whole — PerformanceConfigConvHipConv derives from PerfConfig directly (not PerfConfigBase) and serializes/deserializes the descriptor verbatim, the same way PerfConfigBaseCK handles CK kernel ids. Deserialize clears the resolved index so IsValid() re-resolves the new descriptor; operator== compares descriptors (the identity) rather than the per-build index cache. Adds a CPU-only gtest (conv_hipconv_perf_config.cpp) covering descriptor round-trip and reload through a DbRecord — the path feat(miopen): Cross-architecture AI heuristics for Navis (immediate mode) #9680's picker-replay tests did not exercise.
  • fix: drop ConvHipConv from the perf-config bundle — once the read is fixed the ConvHipConv model becomes reachable, and on the regression set it picks direct[128,1,16,16] almost everywhere (1.2–1.4x slower than the solver default on large shapes, ~0.89x OOTB geomean). Its section is removed from lgbm_pcfg.bin and its golden vectors until the model is retrained and validated on full descriptors. Data-only change: the other 13 solver sections are byte-identical. The picker then abstains for ConvHipConv and the solver uses its own default.

Test plan

  • CPU gtest CPU_ConvHipConvPerfConfig_NONE.* (descriptor round-trip incl. the comma case; reload through DbRecord)
  • Reproduced the bug: the old PerfConfigBase reader truncates direct[tile_size_k=128,...] -> direct[tile_size_k=128
  • Verified the fix round-trips the real struct; conv_hipconv.cpp (stub + MIOPEN_USE_HIPCONV=1), solver.cpp, and the new test compile clean
  • Bundle integrity: 14->13 solvers, other sections byte-identical, golden is a subset of the bundle
  • gfx950 hardware: GPU_UnitTestConvSolverConvHipConv* (search + solution), and a perf-db tune→reload cycle confirming the stored descriptor is not truncated and reloads without an "obsolete or corrupt" warning

Notes

  • No perf-db migration: ConvHipConv ships no db rows yet.
  • Follow-up (out of scope): retrain + re-validate the ConvHipConv pcfg model on full descriptors before re-adding it to the bundle; the separate ASM/CK RGB-layer pick-quality issue; a GPU-gated e2e check that every tunable solver's picked config resolves at runtime.

jdcampbe and others added 2 commits October 8, 2026 09:38
PerformanceConfigConvHipConv serializes its kernel descriptor
(family[field=value,...]). #9680 switched it to this string form but kept
it on PerfConfigBase, whose field reader splits on commas, so a
multi-field descriptor such as
direct[tile_size_k=128,tile_size_n=1,...] was truncated to
direct[tile_size_k=128 on read. The config then matched no kernel, the
perf-config picker skipped every ConvHipConv candidate, and a tuned
config saved to the perf db could not be reloaded.

Serialize/Deserialize the descriptor verbatim, the way PerfConfigBaseCK
handles CK kernel ids, by deriving from PerfConfig directly. Deserialize
clears the resolved index so IsValid() re-resolves the new descriptor, and
operator== compares descriptors (the identity) rather than the per-build
index cache.

Add CPU-only tests covering descriptor round-trip and reload through a
DbRecord, which the picker-replay tests in #9680 did not exercise.

Co-Authored-By: Claude <noreply@anthropic.com>
With the descriptor read fixed, the ConvHipConv perf-config model becomes
reachable, and on the regression set it picks direct[128,1,16,16] almost
everywhere: 1.2-1.4x slower than the solver default on large shapes
(~0.89x OOTB geomean), helping only some small shapes. Remove its section
from lgbm_pcfg.bin and its golden vectors until the model is retrained and
validated on full descriptors. The picker then abstains for ConvHipConv
and the solver falls back to its own default.

Data-only change: the other 13 solver sections are byte-identical.

Co-Authored-By: Claude <noreply@anthropic.com>
@jdcampbe
jdcampbe requested a review from a team as a code owner October 8, 2026 17:44
@therock-pr-bot

therock-pr-bot Bot commented Oct 8, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

therock-pr-bot Bot commented Oct 8, 2026

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

static void Visit(Self&& self, F f)
void Serialize(std::ostream& stream) const override
{
stream.write(descriptor.c_str(), static_cast<std::streamsize>(descriptor.length()));

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.

Suggested change
stream.write(descriptor.c_str(), static_cast<std::streamsize>(descriptor.length()));
stream << descriptor;

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.

[SUGGESTION] The test file is conv_hipconv_perf_config.cpp with suite CPU_ConvHipConvPerfConfig_NONE. The convention for perf-config tests is perf_config_<Solver>.cpp with suites named CPU_PerfConfig_<Solver>_NONE.

for(const auto& desc : Descriptors())
{
// A comma in the stored value is what the field reader split on. ':' and ';'
// are DbRecord's own id/entry separators, which a value must not contain.

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.

@jdcampbe

jdcampbe commented Oct 8, 2026

Copy link
Copy Markdown
Contributor Author

This work is covered by #13213.

@jdcampbe jdcampbe closed this Oct 8, 2026
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