[jdcampbe]

Motivation

MIOpen ships per-arch find/perf databases under src/kernels/, but no gfx950 SystemDB was packaged. Without it, gfx950 runs fall back to runtime tuning on every invocation. MI350X and MI355X both expose 256 CUs, so a single gfx950100 basename serves both per Handle::GetDbBasename() (hex(256) = 0x100, and the _<cu> suffix is omitted for >64 CU).

Technical Details

• Adds projects/miopen/src/kernels/gfx950100.HIP.fdb.txt.bz2 and projects/miopen/src/kernels/gfx950100.db.txt.bz2.
• Adds gfx950 to ALL_GPU_DATABASES in projects/miopen/CMakeLists.txt so install_gpu_databases() picks up the new bz2 files via the existing GLOB pattern.
• No code changes outside the CMake list entry; basename selection, loader, and install paths are unchanged.

Test Plan

• dbsync-test StaticFDBSync/5 round-trip on the staged find DB.
• Confirm install_gpu_databases() packages both gfx950100.*.bz2 blobs in the install tree.
• Smoke MIOpenDriver convolution runs on MI350X / MI355X to confirm the DBs are loaded (no fallback to runtime tuning for shapes present
  in the DB).

Test Result

dbsync-test StaticFDBSync/5 passes on the staged find DB. Runtime smoke on MI350X/MI355X pending; will update with results.

Risk Assessment

Risk level: Low

• Impacted components: projects/miopen/src/kernels/gfx950100.HIP.fdb.txt.bz2, projects/miopen/src/kernels/gfx950100.db.txt.bz2,
  and the ALL_GPU_DATABASES list in projects/miopen/CMakeLists.txt. No source / kernel / solver code is touched. Other archs (gfx942,
  gfx90a, etc.) are unaffected.
  gfx90a, etc.) are unaffected.
• Potential side effects:
  • Functional: gfx950 (MI350X/MI355X) now finds tuned entries instead of falling back to runtime search — strictly additive for those
    archs.
  • Performance: removes per-run tuning cost on covered shapes; no hot-path code change.
  • ABI/API: none.
• Mitigation steps:
  • Change is data + one CMake list entry; trivially revertible by dropping the two bz2 files and the list entry.
  • Find DB was validated via the existing dbsync-test round-trip before staging.
  • No other arch's DB or install logic is affected.

Submission Checklist

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

[yassinsolim]

I don’t think this is ready as-is. The install mechanics and gfx950100 basename look correct, but the DB content is causing performance regressions in focused testing.

I tested the latest PR head (18d7536) on MI350X / gfx950-mi350x with an A/B where the only variable was the system DB path: current staged DB/model tree vs the PR DB overlay. Same current MIOpen/CK libs, same node, same shapes, fresh user DB/cache per run.

Focused examples on MI350X:

bf16 512x256x256 -> 256: ~0.39 ms baseline → ~1.10 ms with the PR DB
bf16 512x1024x1024 -> 256: ~5.92 ms baseline → ~18.39 ms with the PR DB
fp32 256x1024x1024 -> 256: ~12.8 ms baseline → ~51.5 ms with the PR DB
In the focused logs, MIOpen loads matching records from the new gfx950 DB for these cases. So the change is not only removing runtime tuning cost; it can also steer covered shapes to slower recorded entries.

I only have a reliable A/B result on MI350X so far, but since this DB is intended to cover both MI350X and MI355X via the same gfx950100 basename, I think the MI350X regressions are enough to hold this until the entries are refreshed. The CMake/install part looks fine to me; the issue is DB content quality.

[yassinsolim]

Re-ran my A/B on the new head (ce73a60) and this is a lot better, the stuff I flagged before is gone. Couple of new ones to sort out, but we're close.

Setup was the same idea as last time: MI355X, only thing changing is the gfx950 DB. Baseline is 261cdb8 (this branch's libs with the DB taken out, so MIOpen just finds normally), candidate is ce73a60 (same libs + your DB). Same image, same node, same 172 shapes, iter=10, fresh user db/cache each run. Ratio = candidate/baseline, so >1 means the DB entry is slower than what a plain find lands on.

Across the 172 shapes the median is basically 1.0 (0.997) and 67 of them actually got faster. Only 6 are over 1.10x now, vs 42 on the old b84b2a6 DB (which also had that 18x case). So the cleanup clearly worked.

The ones I called out before are all good now:

bf16 512x256x256->256     0.69x   (was ~2.8x)
bf16 512x1024x1024->256   0.71x   (was ~3.1x)
fp32 256x1024x1024->256   0.90x   (was ~4x)

and the 9 MI355 residuals from last round are all 0.74–1.03x now (were 1.06–1.67x).

What's left, these are real DB hits (instant find, so it's the DB picking these, not a search):

bf16 508x2048x2048->256    5.66x    33 -> 186 ms   DB picks ConvHipImplicitGemmGroupFwdXdlops, GemmFwdRest is ~6x faster
fp32 1024x1448x1448->256   1.33x    94 -> 125 ms   DB picks GemmFwdRest, a find lands on something faster
fp32 1024x1024x1024->256   1.05x    48 -> 50 ms    minor

The bf16 508x2048x2048->256 one is the only thing I'd block on, it's a covered shape getting steered to a kernel ~6x slower than it should. Keys if it helps:

508-2048-2048-3x3-256-2048-2048-1-1x1-1x1-1x1-0-NCHW-BF16-F
2x508x2048x2048x1x3x3x1x256x1x1x1x0x1x1x0x1x1x0x0x1xNCHWxBF16xF

Can you re-tune those couple of entries (or just drop them so they fall back to find)? The 1024x1448x1448 fp32 one would be nice to grab too. Everything else looks good to me.

Two caveats so the numbers are in context: I ran this on the branch's current libs, which are ~10 days behind develop, so once you rebase I'd like to do one more quick pass before it merges (CK moves around enough that a couple shapes can shift). And the baseline is a full find, i.e. close to best-case, so these ratios are "DB entry vs the best you can do" rather than vs some weak default, the two I flagged are genuinely bad recordings, not noise.

[yassinsolim]

Re-ran broader validation on the latest f8a1dd2 head. The original ByteDance blocker looks fixed, but I found two reproducible tiny-shape regressions that look related to the DB not including/selecting ConvDirectNaiveConvFwd.

Method:

• candidate = f8a1dd2 image with system gfx950100 DB enabled
• baseline = same image/libs with system find DB disabled, so this isolates the PR DB effect
• repeated the two worst cases 3x

Regressions:

1. MIOpenDriver convfp16 -n 1 -c 1542 -H 14 -W 14 -k 1542 -y 3 -x 3 -p 1 -q 1 -u 2 -v 2 -l 1 -j 1 -g 1542 -m conv -F 1

   • DB: GemmFwdRest, geo ~0.0269 ms
   • no system DB: ConvDirectNaiveConvFwd, geo ~0.00452 ms
   • DB is ~5.9x slower

2. MIOpenDriver convfp16 -n 1 -c 1632 -H 12 -W 12 -k 1632 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -g 1632 -m conv -F 1

   • DB: GemmFwdRest, geo ~0.0301 ms
   • no system DB: ConvDirectNaiveConvFwd, geo ~0.00638 ms
   • DB is ~4.7x slower

If excluding naive from gfx950 SystemDB is intentional policy, please confirm that these tiny-shape regressions are acceptable. Otherwise I think these entries should be dropped or retuned so they fall back to the faster no-DB path.