Skip to content

Problem cache follow-on: pluggable backends, layered caches, compile-options API + offline aggregation/driver tooling - #47

Draft
danieyan-amd wants to merge 31 commits into
feature/problem-cache-schema-extensionfrom
feature/problem-cache-followon
Draft

Problem cache follow-on: pluggable backends, layered caches, compile-options API + offline aggregation/driver tooling#47
danieyan-amd wants to merge 31 commits into
feature/problem-cache-schema-extensionfrom
feature/problem-cache-followon

Conversation

@danieyan-amd

Copy link
Copy Markdown
Owner

DRAFT — preview, stacked on ROCm#4835. Opened inside the fork against the feature/problem-cache-schema-extension (ROCm#4835) branch, so the diff shows only the follow-on work. This will be retargeted to ROCm/AMDMIGraphX:develop and opened for review once ROCm#4835 merges — please don't merge this fork-internal draft.

Builds on ROCm#4835 (device-keyed problem cache) to add the full production tooling around it. This is the second and final problem-cache PR — it turns the engine from ROCm#4835 into an end-to-end workflow for collecting, merging, shipping, and loading tuned kernel solutions. Depends on ROCm#4835 (must merge first).

What's included

  • Pluggable backends (te.py): type-erased problem_cache_backend + a JSON concrete backend (json_problem_cache).
  • SQLite backend: sqlite_problem_cache as a second backend (vendored SQLite amalgamation).
  • Explicit cache path: compile_options.problem_cache_path + a problem_cache::load(path) overload.
  • Offline aggregator: merge / validate / convert (device-keyed merge, conflict policy, dedup, reporting).
  • Driver subcommands: aggregate_cache, validate_cache, convert_cache for offline cache management.
  • Multi-cache priority list: layered loading (app > local > shipped) — read all in priority order, write to the writable layer.
  • GFX-arch remapping: fold minor GFX variants to a canonical arch (e.g. gfx1151 -> gfx1150) to dedup across compatible GPUs (--remap-gfx).

Testing

  • Builds clean: migraphx_gpu + migraphx-driver link (only pre-existing warnings).
  • 41 unit tests pass, 0 failures: problem_cache_backend (4), sqlite_problem_cache (5), problem_cache_path_override (2), problem_cache_aggregator (29), test_sqlite (1).
  • Driver verified on real data (51 gfx1100 cache files): aggregate_cache (first-wins) -> 672 entries; validate_cache -> valid; convert_cache json->json and json->sqlite (672); --remap-gfx works; error paths report cleanly (no crash) on bad options / conflicts.

Notes

Single follow-on PR (all features in one), stacked on ROCm#4835. Includes two small reconciliation/hardening commits (adapt to ROCm#4835's set_device_key(context) API; graceful driver error reporting) — squashable into their feature commits before the real upstream PR.

dependabot Bot and others added 14 commits July 23, 2026 14:40
Macro builders were recently added to better manage composed ops when parsing from outside libraries. This extends the torch kit with a number of missing ops required to fully migrate torch_migraphx to use this builder framework.
Adds gfx115 to hipblaslt_supported_impl() so gfx1150/1151/1152/1153 (Strix Halo/Point, RDNA3.5) use the hipBLASLt GEMM path.
…nd update tests (ROCm#5090)

Rewrites the flash decoding kernel 2 recombination step in find_flash_decoding to use the exp-normalize form:

`O = sum(O' * exp(LSE - max)) / sum(exp(LSE - max))`

instead of normalizing weights first, then scaling and summing partial outputs. The result is mathematically equivalent but produces IR that fuses more cleanly downstream (e.g. with rewrite_broadcast in a follow-up PR).
@danieyan-amd
danieyan-amd force-pushed the feature/problem-cache-followon branch from f33193b to 948355f Compare July 29, 2026 18:29
@danieyan-amd danieyan-amd changed the title Problem cache follow-on: pluggable backends, offline aggregation + driver tooling, layered caches, GFX remapping Problem cache follow-on: pluggable backends, layered caches, compile-options API + offline aggregation/driver tooling Jul 29, 2026
@danieyan-amd
danieyan-amd force-pushed the feature/problem-cache-followon branch from 948355f to caa0c6d Compare July 30, 2026 05:17
@danieyan-amd
danieyan-amd force-pushed the feature/problem-cache-followon branch 2 times, most recently from 3798b85 to 0ba5b7b Compare August 5, 2026 21:32
@danieyan-amd
danieyan-amd force-pushed the feature/problem-cache-followon branch from 8c119b5 to 6045418 Compare August 7, 2026 00:31
danieyan-amd added 3 commits August 6, 2026 22:53
Introduce a type-erased problem_cache_backend (pluggable storage) with JSON and SQLite implementations; the runtime problem_cache routes has/insert/mark/get/load/save through it and canonicalizes keys so JSON round-trips match shipped caches. Move cache_device_key to its own header and route gpu/ DSL generation via generate.py.

Signed-off-by: danieyan-amd <daniel.anieyan@amd.com>
Load an ordered list of read-only caches (first hit wins) alongside a writable cache; compile_ops and the gemm paths query them during tuning.

Signed-off-by: danieyan-amd <daniel.anieyan@amd.com>
Read problem_cache_files from the GPU backend options and load them into the context; no dedicated API or environment variable is required.

Signed-off-by: danieyan-amd <daniel.anieyan@amd.com>
@danieyan-amd
danieyan-amd force-pushed the feature/problem-cache-followon branch from 6045418 to 2911a33 Compare August 7, 2026 13:24
Addresses review feedback: the multi-cache priority search lived in the
context (read_only_caches + find_in_problem_caches). Fold it into
problem_cache so has() and get() search the read-only layers (highest
priority first) then the writable cache. The context now just delegates.
Behaviour is unchanged: first hit wins; a single file is writable and
multiple files are a read-only priority list.

Signed-off-by: danieyan-amd <daniel.anieyan@amd.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.