[amd-ahyttine]

JIRA ID : SILOTIGER-1145

Motivation

On gfx942/gfx950, the AI-based CandidateSelectionModel heuristics are invoked during solver
selection to rank CK kernel configurations. For workloads that encounter the same convolution shape
more than once within a process (e.g. training loops), or the same solver across different shapes,
the model runs in full every time. This overhead can dominate kernel execution time for small shapes.
This PR adds two lightweight in-process RAM caches to eliminate redundant inference.

Risk Assessment

Low risk. All changes are confined to ai_candidate_selection.cpp and its header, plus the
corresponding test file. No public APIs, ABI, kernel implementations, build system, or database
contents are modified. The two caches are in-process only and are bypassed entirely when
MIOPEN_ENABLE_AI_KERNEL_TUNING is off. On a cache miss the code path is identical to the
pre-PR behaviour, so the only observable difference on a hit is a faster return with the same
result. The result cache stores and returns the ranked CandidateSelectionResult verbatim;
the embedding cache is guarded against collisions by using exact float bit-pattern keys. Both
caches are protected by per-instance mutexes. The risk of regression is limited to workloads
that call ModelSelectBestCandidate for the same (arch, solver, problem, use_split_k)
more than once within a process, which is also where the benefit is observed.

Technical Details

Two caches are added inside ModelSelectBestCandidate() in ai_candidate_selection.cpp:

1. Result cache — caches the full CandidateSelectionResult (ranked kernel indices + split_k
   values) keyed on (arch, solver, problem_description, use_split_k). A hit bypasses all model
   inference entirely, mirroring the existing TunaNet RAM cache in ai_heuristics.cpp.

2. Kernel-config embedding cache — the kernel-config encoder output depends only on the kernel
   parameters, not on the input shape. Embeddings are cached at individual kernel-row granularity: on
   each call only the rows not already in cache are encoded (in a single batched fdeep call), and the
   results are stored per row. This allows partial reuse across different shapes that share any subset
   of valid CK kernel instantiations, which is the common case when layers differ in spatial dimensions
   but share the same solver.

ModelSelectBestCandidate() gains a const conv::ProblemDescription& parameter (needed as the
result cache key); both call sites in ai_conv_nd_kernel_tuning_utils.hpp are updated accordingly.

Test Plan

1. Run the existing + new unit tests:
   ./bin/test_conv_ai_candidate_selection_model

Three new test cases are added to GPU_CandidateSelection_FP32:

• CandidateSelectionRamCache_Test — calls ModelSelectBestCandidate twice with identical
  arguments, asserts identical results, logs timing of both calls.
• KernelEmbeddingCache_Test — encodes a set of kernels, then encodes a strict subset of those
  same kernels, and asserts the overlapping rows produce byte-identical embeddings, confirming
  per-row cache correctness.
• CandidateSelectionRamCacheIsolation_Test — populates use_split_k=false and use_split_k=true
  cache entries independently and verifies they do not contaminate each other.

2. Run a simple convolutional network training/inference in Pytorch using IGEMM kernels and confirm operation.

Test Result

1. All 38 tests in GPU_CandidateSelection_FP32 pass. Cache-hit log lines confirmed at log level 6
   (MIOPEN_LOG_LEVEL=6) on the second call in each relevant test.

2. Timings from a 2-layer convolutional network run on MI300 (docker image built from each branch).
   First-call overhead includes heuristic inference + kernel compilation. Subsequent calls are
   served from MIOpen's invoker cache (no heuristics involved).

Layer	Direction	Call	Develop (ms)	This branch (ms)	Delta
features.0 (Conv2d)	Forward	1st	5009.0	4743.7	−265 ms (−5.3%)
features.3 (Conv2d)	Forward	1st	20.2	19.7	−0.5 ms (−2.5%)
features.0 (Conv2d)	Forward	2nd+	0.78	0.80	≈0 (invoker cache)
features.3 (Conv2d)	Forward	2nd+	0.67	0.62	≈0 (invoker cache)
features.3 (Conv2d)	Backward	1st	114.5	109.3	−5.2 ms (−4.5%)
features.0 (Conv2d)	Backward	1st	0.04	0.04	≈0
features.3 (Conv2d)	Backward	2nd+	1.58	1.40	≈0 (invoker cache)

The 2-5% speedup on the first forward and backwards calls to features.3 corresponds to the kernel-config
embedding cache hit logged as:

Kernel-config embedding cache: 5/69 kernels hit, encoding 64 new kernels

The 5 kernel rows shared with a previous solver call are not re-encoded. The improvement scales
with the number of shared kernel rows across solvers and shapes; larger networks with more
repeated solver calls will see proportionally greater benefit.

Submission Checklist

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

[JonathanLichtnerAMD]

Can we forward declare this instead of adding the include?

[amd-bartgips]

Fixed. #include <miopen/conv/problem_description.hpp> has been removed from ai_candidate_selection.hpp and replaced with a forward declaration of conv::ProblemDescription. The full include is now in ai_candidate_selection.cpp only.

[JonathanLichtnerAMD]

Claude found this one:

KernelEmbeddingCache_Test does not test the cache. The test calls model.EncodeKernelConfigs(...) directly, which is a thin wrapper around EncodeKernelConfigsWithFdeep and never
touches KernelEmbeddingCache. The cache is only invoked from GetOrComputeKernelEmbeddings (anonymous namespace, ai_candidate_selection.cpp:106-168). The test would pass even if the cache
code were deleted. Fix by exercising via ModelSelectBestCandidate with overlapping kernel sets, or by exposing the helper.

[amd-bartgips]

Fixed.
Claude tells me:

The test previously called model.EncodeKernelConfigs() directly, which bypasses GetOrComputeKernelEmbeddings entirely. It has been rewritten to exercise the cache via ModelSelectBestCandidate: two calls with identical features and kernel params but different ProblemDescription objects (different B-RAM keys, same embedding cache keys). Because the scorer receives identical Tower 1 output and identical cached Tower 2 embeddings, the rankings must be deterministically the same — a failure would indicate the cache returned a corrupted or wrong embedding.

[JonathanLichtnerAMD]

Is there a chance that we get collisions here? From claude:

HashKernelRow collision risk (ai_candidate_selection.cpp:97-104). The size_t hash IS the key; there is no post-hash equality check on the float rows. With small-integer-valued
encodings from EncodeKernelParams, std::hash + boost-XOR-shift produces structured output that elevates collision rate above the 64-bit theoretical floor. A collision returns the
wrong embedding silently and SelectBestCandidateIndices then ranks against it. Fix: build the key from a delimited std::to_string (or hex-encoded bit pattern) of the row.

[amd-bartgips]

Fixed.
I recall Claude also flagging something similar but deemed the likelihood of this happening to be extremely low. Regardless, there's no reason to not properly fix this.

The hash-based key has been replaced with an exact bit-pattern key: each float in the row is memcpy-ed to a uint32_t and appended as hex to the key string, prefixed by arch|solver|row_size. This is collision-free by construction. The function has been renamed KernelRowKey to reflect that it is no longer a hash.

[JonathanLichtnerAMD]

We also need to add a Risk Assessment section to this PR message (although I believe we can drop that once we merge).

[amd-bartgips]

We also need to add a Risk Assessment section to this PR message (although I believe we can drop that once we merge).

Thanks Jonathan, I responded to all your comments and added the RA section to the PR text.

[github-actions]

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

• Add new commits
• Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

[github-actions]

This pull request has been automatically closed due to inactivity (30 days with no updates).

If you'd like to continue working on this, feel free to reopen the PR or create a new one.

[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.

[amd-bartgips]

I think we can probably close this again, as it seems to be superseded by #12947

[github-actions]

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

• Add new commits
• Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

[amd-bartgips]

Superseded by #12947