[JonathanLichtnerAMD]

JIRA ID: ROCM-27477

Motivation

The AI convolution heuristics in MIOpen cache loaded models in process-global static maps (GetModel and GetFdeepModel in ai_heuristics.cpp). When convolutions run concurrently from multiple threads on the same device (for example, per-thread MIOpen handles each running a convolution), those callers race on the same cache key during first population. This unguarded find-then-insert was a data race that could corrupt the map and produce a use-after-free deep inside fdeep predict, segfaulting the process (SIGSEGV/SIGABRT). The crash was first surfaced via ConvHipImplicitGemmGroupFwdXdlops's HeuristicInit.

This is a cherry-pick / backport of the already-merged PR #9131 onto the release/therock-7.14 branch.

Technical Details

projects/miopen/src/conv/heuristics/ai_heuristics.cpp:

• Added #include <mutex>.
• GetModel(arch, solver): added a static std::mutex models_mutex and a std::lock_guard held across the find + insert so cache population is serialized. Moved the models.find lookup to occur after the lock is acquired (and after the gfx950 to gfx942 arch remap). Collapsed the if/else into a single trailing return it->second.
• GetFdeepModel(path, key): changed the cache from std::map<std::string, std::unique_ptr<fdeep::model>> to std::map<std::string, std::shared_ptr<fdeep::model>> and changed the return type from const fdeep::model& to std::shared_ptr<fdeep::model>. Added a static std::mutex models_mutex guarding find + insert. Returning by shared_ptr (rather than a reference into the map) keeps the model alive in the caller independently of the cache, so it stays valid even if the entry were ever replaced or evicted, mirroring GetModel.
• Updated the two call sites in EncodeInputFeaturesWithFdeep and EncodeKernelConfigsWithFdeep to use the returned shared_ptr (->predict(...), ->predict_multi(...)) instead of a reference into the map.

Trade-off: serializing the cache with a mutex adds contention only on the first-population path; subsequent lookups are cheap hash-map reads under a short-held lock. Returning by shared_ptr adds a refcount bump per call, which is negligible relative to model inference cost.

projects/miopen/test/gtest/conv_ai_heuristics_threading.cpp (new, 173 lines): a gtest that hammers GetFdeepModel from 32 threads over 8 iterations on a cold cache, each thread owning its own CandidateSelectionModel (mirroring per-thread handles), lined up with a std::latch to maximize the first-population overlap. Success requires no crash, no thrown exception, and identical non-empty encodings across all threads (a corrupted/half-built model would diverge or throw). Gated to Instinct/CDNA devices (candidate-selection KTN models ship only for gfx942/gfx950) and skips when model files are absent. Compiled under MIOPEN_ENABLE_AI_KERNEL_TUNING.

Risk Assessment

Risk Level: 🟢 Low

The change is a small, contained thread-safety fix (39 lines changed in one source file plus a new test file). It adds locking to two cache accessors and shifts one cache from unique_ptr to shared_ptr; it does not alter heuristic logic, model outputs, or any public API. The only behavioral change is correctness under concurrency.

Impacted Components

• MIOpen AI convolution heuristics: projects/miopen/src/conv/heuristics/ai_heuristics.cpp (GetModel, GetFdeepModel, EncodeInputFeaturesWithFdeep, EncodeKernelConfigsWithFdeep).
• MIOpen gtest suite: new projects/miopen/test/gtest/conv_ai_heuristics_threading.cpp.
• Indirectly: any convolution path that invokes the AI candidate-selection / model heuristics, especially in multi-threaded host usage.

Potential Side Effects

• Slight added lock contention on first-population of each cache key; steady-state impact is negligible (short-held lock on a hash-map read).
• GetFdeepModel now returns a shared_ptr rather than a reference; the two internal call sites were updated accordingly. The function is internal to the candidate_selection namespace, so no external callers depend on the old signature.
• No change to numerical results or selected kernels.

Mitigation Steps

• Added a dedicated concurrency regression test (ConcurrentGetFdeepModel) that reproduces the original race on a cold cache and would also trip ThreadSanitizer.
• Change is self-contained in a single commit and trivially revertable.
• This is a backport of code already merged and reviewed upstream in PR fix(miopen): guard AI heuristic fdeep model caches with a mutex #9131.

Test Plan

• Build the MIOpen gtest binary and run GPU_AIHeuristicsThreadSafety_FP32.ConcurrentGetFdeepModel (parameterized over gfx942 and gfx950) on an Instinct/CDNA device with candidate-selection models shipped. The test spawns 32 threads x 8 iterations against a cold cache.
• Optionally run under ThreadSanitizer for an order-independent data-race check.
• The test self-skips on non-Instinct devices and when model files are absent, so it is safe to include in broad CI runs.

Test Result

Upstream PR #9131 is merged and reviewed: before the fix the concurrent burst crashes (SIGSEGV/SIGABRT) or trips ThreadSanitizer; with it, all threads return identical, non-empty encodings with no crash or exception. This cherry-pick applies as a single commit with no additional modifications. Final validation is expected via the MIOpen gtest suite on the release/therock-7.14 CI targets.

[marbre]

Looks go from the perspective that it cherry-picks a commit that already made it into develop. One of the miopen shards failed tests (I retriggered) but I assume this is flaky. Build tests for gfx950 are unrelated and build passed on the original PR (so did tests). The risk here however is that the cherry-picked commit did not yet propagate into TheRock (this needs another submodule bump) and therefore we didn't had a nightly build yet.

Approving for DevOps though as we could catch regressions either on the submodule bump or on the RC build.

[marbre]

gfx94X went through, gfx950 went through on the original PR. The prior build failure was unrelated and can therefore be ignored.