Problem cache follow-on: pluggable backends, layered caches, compile-options API + offline aggregation/driver tooling - #47
Draft
danieyan-amd wants to merge 82 commits into
Conversation
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
force-pushed
the
feature/problem-cache-followon
branch
from
July 29, 2026 18:29
f33193b to
948355f
Compare
danieyan-amd
changed the title
danieyan-amd
force-pushed
the
feature/problem-cache-followon
branch
from
July 30, 2026 05:17
948355f to
caa0c6d
Compare
danieyan-amd
force-pushed
the
feature/problem-cache-followon
branch
2 times, most recently
from
August 5, 2026 21:32
3798b85 to
0ba5b7b
Compare
Some rocMLIR tuning configurations require splitting fused convolution/GEMM and pointwise modules. The resulting pointwise partition may contain unsupported shape operations, causing compilation failures.
Expose separate read-only problem_cache_files (shipped gpuep/ISV caches, never written back) and writable_problem_cache_files (the developer cache that new tuning solutions save back to). Addresses review feedback to offer distinct read-only and read/write cache configuration. Signed-off-by: danieyan-amd <[email protected]>
…view-fixes # Conflicts: # src/targets/gpu/target.cpp
Cover the two-tier scenarios end to end: writable-only and no-cache configs, a byte-for-byte check that read-only caches are never rewritten on save, and backend-option parsing that keeps the read-only and writable options distinct. Signed-off-by: danieyan-amd <[email protected]>
danieyan-amd
force-pushed
the
feature/problem-cache-followon
branch
from
August 20, 2026 19:08
b8c059d to
fabcd93
Compare
Store weights in fp32 registers and accumulate in fp32 to improve accuracy for fp16 input. Cast back to output type at the end.
…#5112) * Dynamic slice from symbolic slice in other branch * Test and comment cleanup * Cleanup normalize_attributes * Fix licensing * Simplify comment * Add fold_min and fold_max to dyn_slice compute_shape() * Add throw for if extent is always negative * use dim_like and 1_arg slice hold * Add normalize_attributes tests * cleanup * Remove unused dim_like functions * cleanup * more cleanup * Add more tests * Cleanup 3 * licensing * Apply suggestions from code review Co-authored-by: Copilot Autofix powered by AI <[email protected]> Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * Use exact match for dyn_slice ref tests * Add min() and max() same symbol rewrite rules * Allow creating symbolic literal from number or string * Fix test * Move eval_expr test from slice to dyn_slice test * Rename fold_min and fold_max to resolve_min and resolve_max --------- Co-authored-by: Copilot Autofix powered by AI <[email protected]> Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
…ve no undefined symbols (ROCm#5154)
Make gemm_default_solution and hip_gemm_default_solution take a mutable context so they use the non-const get_problem_cache(), and remove the now-unused const overload. Fold problem_cache::load(std::string) into load(read_only, writable) and update the path-override tests to the two-list form. Signed-off-by: danieyan-amd <[email protected]>
Replace the single writable backend with parallel writable_backends and save_paths vectors: load every writable file, save each back, and seed an in-memory json cache when none is configured. Rename make_backend to make_problem_cache_backend and drop the auto_save flag (save() is a no-op without a save path). Rename the GPU target options so problem_cache_files is the writable default and read_only_problem_cache_files holds the system-level caches. Signed-off-by: danieyan-amd <[email protected]>
Signed-off-by: danieyan-amd <[email protected]>
The env var was removed in favor of the GPU backend option; drop the leftover entry in the dev env-vars doc. Signed-off-by: danieyan-amd <[email protected]>
danieyan-amd
force-pushed
the
feature/problem-cache-followon
branch
from
August 25, 2026 21:11
6ecb5ab to
c79a34c
Compare
Search the writable and read-only backend vectors with std::find_if instead of raw range-for loops. Behavior is unchanged (writable caches first, first hit wins). Signed-off-by: danieyan-amd <[email protected]>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
19 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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 toROCm/AMDMIGraphX:developand 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
problem_cache_backend+ a JSON concrete backend (json_problem_cache).sqlite_problem_cacheas a second backend (vendored SQLite amalgamation).compile_options.problem_cache_path+ aproblem_cache::load(path)overload.merge/validate/convert(device-keyed merge, conflict policy, dedup, reporting).aggregate_cache,validate_cache,convert_cachefor offline cache management.--remap-gfx).Testing
migraphx_gpu+migraphx-driverlink (only pre-existing warnings).problem_cache_backend(4),sqlite_problem_cache(5),problem_cache_path_override(2),problem_cache_aggregator(29),test_sqlite(1).aggregate_cache(first-wins) -> 672 entries;validate_cache-> valid;convert_cachejson->json and json->sqlite (672);--remap-gfxworks; 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.