[jdcampbe]

Motivation

gfx942 encompasses multiple cards. The current system dbs don't cover all of the cards due to inexact ("closest compute-unit") fallback silently selected nothing, so MIOpen ran with no system tuning data and logged File is unreadable: ".../miopen_system_db/gfx942_64.HIP.fdb.txt". This PR fixes the root-cause bug that left the fallback's candidate list empty, and adds a small, scoped redirect so such parts fall back to the appropriate sibling gfx942 database (gfx942130) instead of the nearer-by-count but less suitable gfx942e4.
JIRA ID : ALMIOPEN-2121

Technical Details

System DBs are named an exact base naming system. When the exact basename file is absent, MIOpen runs an inexact search: scan the DB directory and pick the file for the same arch (DbId) with the numerically closest CU count.

Root-cause bug. The inexact search built its candidate list with fs::path::extension():

if(fs::is_regular_file(filepath) && filepath.extension() == path_suffix + ".fdb.txt")

std::filesystem::path::extension() returns only the final component after the last dot — .txt for gfx942e4.HIP.fdb.txt — so it can never equal the multi-dot .fdb.txt / .db.txt. On disk-based text DBs (every shipped DB) the candidate list was always empty, the inexact search returned an empty path, and no fallback DB was loaded. This was invisible on devices that ship their own DB because they short-circuit on the exact-match fs::exists() path, which never touches this filter.

Fix 1 — candidate collection. Match against the full filename with EndsWith(...) instead of extension(), in all four inexact loops:

• src/find_db.cpp:163 (find-db, on-disk) and the embedded find-db loop (already used EndsWith, left as is at src/find_db.cpp:83).
• src/include/miopen/execution_context.hpp:230 (perf-db, on-disk) and :153 (perf-db, embedded); added #include <miopen/stringutils.hpp>.

Fix 2 — scoped sibling redirect. Add GetSysDbSelectionCu() in src/include/miopen/handle.hpp:

inline int GetSysDbSelectionCu(const std::string& db_id, int real_cu_count)
{
    if(db_id == "gfx942" && real_cu_count <= 80)
        return 304;
    return real_cu_count;
}

Intentionally not changed. User DBs (*.udb.txt / *.ufdb.txt) keep the device's own basename and have no inexact fallback, so per-device tuned data is still read and written separately. The exact-match path is untouched — devices with their own DB short-circuit before any modified code runs. GetSysDbSelectionCu returns the real CU count unchanged for every other arch and for gfx942 parts above the threshold. The kdb path (arch-only fallback) is unaffected.

Test Plan

Build a multi-arch image from this branch and run a small smoke set (conv fwd/bwd/wrw x fp32/fp16/bf16) with MIOPEN_LOG_LEVEL=7 (to capture DB resolution) and MIOPEN_FIND_MODE=2, across:

• mi308x (no dedicated DB) — the fallback target,
• mi325x/mi300x — sanity,
• a gfx950 — sanity.

For each, inspect the shard log for the find-db / perf-db resolution lines and confirm:

• GOOD markers: Found exact ... database file for exact parts; inexact find database search then FindRecord ... in file ".../<expected>" for the fallback part.
• No File is unreadable on a miopen_system_db path.

Test Result

Part	find-db resolved	perf-db used	system-DB unreadable?	Result
gfx942, low-CU (fallback)	inexact -> gfx942130	gfx942130	none	fallback now works
gfx942, 304 CU (sanity)	exact gfx942130	gfx942130	none	no regression
gfx950 (sanity)	exact gfx950100	gfx950100	none	no regression

The fallback part now logs inexact find database search, then FindRecord ... in file ".../gfx942130.HIP.fdb.txt", and the perf-db opens gfx942130.db.txt (the transient gfx942e4 "best candidate" line is superseded by gfx942130 at distance 0 from the biased target; only gfx942130 is actually opened). No File is unreadable on the system DB. Exact-match parts log Found exact ... database file and load their own DBs unchanged. Two additional same-DB sanity runs on a second node each (mi300x/mi325x and gfx950) are still in progress and are expected to match the corresponding sanity rows; will update if anything differs.

Risk Assessment

Risk level: Low

• Impacted components: system-DB path resolution only — src/find_db.cpp, src/include/miopen/execution_context.hpp, src/include/miopen/handle.hpp (3 files, +40 / -11). Affects only the inexact ("closest-CU") fallback branch. Not touched: the exact-match path, user DBs, the kdb path, public API/ABI, kernels, and any non-gfx942 arch.
• Potential side effects:
  • Functional: previously-broken fallback parts now load a usable system DB; low-CU gfx942 parts resolve to gfx942130 rather than gfx942e4 or nothing. No behavior change for any device that already resolved a DB.
  • Performance: negligible — one extra integer comparison per DB lookup, on the cold DB-open path only.
  • ABI/API: none.
• Mitigation steps:
  • Change is small and localized to the inexact branch; the exact-match path that all DB-shipping devices use is provably unentered (short-circuits on fs::exists() before the modified code).
  • Defect and fix are easy to reproduce and verify pre-merge with MIOPEN_LOG_LEVEL=7 (see Test Plan); validated on real gfx942 and gfx950 hardware.
  • Fully revertible — no data files added, renamed, or removed.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[therock-pr-bot]

✅ 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]

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

[cderb]

LGTM