Repository navigation
fix(miopen): guard AI heuristic fdeep model caches with a mutex - #9131
JonathanLichtnerAMD merged 3 commits into
Conversation
GetFdeepModel and GetModel in ai_heuristics.cpp cache models in a static std::map with no synchronization. Under concurrent convolution from multiple threads on the same device (e.g. per-thread MIOpen handles), two callers racing on the same key both insert; the second insert move-assigns over the first's unique_ptr, destroying a model another thread is still using through the returned reference -> use-after-free / SIGSEGV deep in fdeep predict. Serialize find+insert with a static std::mutex, mirroring the existing GetCandidateSelectionModel. GetFdeepModel returns a reference to a map entry that is never overwritten afterward, so it stays valid once the lock is released; GetModel returns its shared_ptr by value under lock. ROCM-27477 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
gfx942 (MI300X) validation — fixes the
|
| Build | Runs | Crashes |
|---|---|---|
Unpatched (develop − fix) |
220 | 3 — 2× SIGSEGV(11), 1× SIGABRT(6) |
| Patched (this PR) | 220 | 0 |
- Crash signals (SIGSEGV + SIGABRT/heap-corruption) match the RCA. Core dumps confirm the fault is inside
libMIOpen.so.1with a corrupted stack — the use-after-free signature (the Release build strips the internalfdeep/GetFdeepModelsymbols, so exact frame names don't resolve, but the crash region and corruption pattern match). - Observed crash rate here was ~1.4%, lower than the originally reported 3–17% — expected, since these runs were spread across 4 GPUs in parallel on a busy shared host, which shifts the timing of the warm-up race. The key point is the clean separation: the crash never occurred with the fix (0/220), and reproduced only without it.
BrianHarrisonAMD
left a comment
There was a problem hiding this comment.
Changes LGTM.
I think adding some tests for model retrieval is worth it.
GetFdeepModel cached fdeep models in a static map of unique_ptr and handed callers a bare `const fdeep::model&` into the map. That is only safe while entries are never overwritten; a future change (cache eviction/replacement) would silently reintroduce a use-after-free. Store the models as shared_ptr and return them by value, mirroring GetModel. The returned model now stays alive for the caller independently of the cache, so the lock only needs to serialize population, not guard every subsequent read. ROCM-27477 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Add a Standard-category gtest that hammers GetFdeepModel from 32 threads, lined up with a std::latch to hit the first-population window. Asserts no crash, no exceptions, and identical non-empty encodings across threads. The work is host-side (fdeep load/predict on CPU), but the test is gated to Instinct/CDNA devices via IsTestSupportedByDevice, since the candidate-selection models only ship for gfx942/gfx950 and it should not run on consumer/APU parts (gfx1151, gfx1103, ...). Parameterized over gfx942/gfx950 and skips when those models are not shipped. Catching the race needs a cold cache at burst start. The gfx950 parameter is reliably cold (no other test uses gfx950 candidate-selection keys); the gfx942 key can be pre-warmed by conv_ai_candidate_selection_model when the whole miopen_gtest binary runs in a single process, though a Standard-only or isolated run keeps it cold. Reproduces the ROCM-27477 use-after-free (SIGSEGV) reliably against the pre-fix code with a cold cache, and passes with the fix. ThreadSanitizer would be the order-independent guard. ROCM-27477 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
b2113ff to
145d7ab
Compare
…ry-pick #9131) (#9149) 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 #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. Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
JIRA ID: ALMIOPEN-2312 ## Motivation `projects/miopen/CHANGELOG.md`'s "MIOpen 3.5.2 for ROCm 7.14.0" section only had 4 entries, but 68 commits touched `projects/miopen` between the `therock-7.13` tag and `release/therock-7.14`. Auditing that commit range (via `git log --oneline therock-7.13..remotes/origin/release/therock-7.14 -- projects/miopen`) found 18 additional user-facing changes — new gfx950/gfx1250 architecture support, SystemDB refreshes, AI-heuristic improvements, several correctness/crash bug fixes, and OpenCL solver removals — that were never documented. This PR backfills those entries so the 7.14.0 changelog accurately reflects what shipped. ## Technical Details Added 24 lines to the existing "MIOpen 3.5.2 for ROCm 7.14.0" section of `projects/miopen/CHANGELOG.md`, merging new bullets into the existing `Added`/`Changed`/`Resolved Issues` sub-sections (no new entries were duplicated or reworded) and adding a new `Removed` sub-section for solver removals. No existing content was altered. Each new bullet was traced to a specific commit + PR: | Entry | Commit | PR | |---|---|---| | Added gfx950 (MI350X/MI355X) SystemDB | `60d8e1e5175` | #8081 | | Added initial gfx1250 architecture support | `d3f057b2f22` | #7587 | | Added JSON-formatted performance logs (`MIOPEN_PERFORMANCE_LOGS`) | `ffbde871260` | #5736 | | Refreshed gfx942 SystemDB | `faf754a9ef6` | #8174 | | Refined 2D AI heuristics (Tunanet/KTN) for gfx942 & gfx950 | `c24ad415884` | #7109 | | Fixed thread-safety issue in AI-heuristic model caches | `cd957402309` | #9149 (cherry-pick of #9131) | | Fixed CK grouped-conv solvers not always registered in host | `c971d54d7ae` | #9105 | | Fixed legacy CK buffer resource descriptor (gfx1250) | `1dee011b8b0` | #8758 | | Removed stale gfx942 SystemDB naive-solver entries | `f42668f6316` | #7028 | | Fixed softmax NaN when `beta=0` | `fb8157c8512` | #8426 | | Fixed BN backward spatial variant 1 vectorized memory access bug | `1351a7e8b10` | #6765 | | Fixed use-after-free of `kernel_id` in `InitInvokerFactoryNCHW` | `394ebffb50d` | #7922 | | Fixed SoftMarginLoss forward reduce-kernel kernarg corruption | `1e65a31570f` | #8065, #8135 | | Fixed build with `MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF` | `e6288c3ca61` | #6800 | | Fixed install RPATH for MIOpenDriver / CK backend libs | `ef8a4cf2f93` | #7485 | | Fixed dangling lambda causing log-level 6 segfault | `190f9791737` | #6953 | | Enabled Windows long-path support for `addkernels` | `1f9d69724cf` | #7387 | | Removed OCL solver `ConvOclDirectFwd1x1` | `126c720cadd` | #6491 | | Removed OCL solver `conv_ocl_dir2Dfwdgen` | `2d499ddf709` | #6633 | | Removed OCL solvers `ConvOclBwdWrW1x1`, `ConvOclBwdWrW2` | `b5d43a33dc6` | #6859 | | Removed OCL solvers `OpTensorLeadingOnes`, `Op5dTensorGeneric` | `8642c83b7de` | #7347 | The remaining ~40 commits in the 68-commit range were deliberately excluded (tests/CI/build/docs/internal refactors/1:1 OCL→HIP swaps with no behavior change, per each commit's own stated scope) and are not represented here. ## Risk Assessment **Risk Level: 🟢 Low** This is a documentation-only change to a single Markdown file. No code, build configuration, tests, or public API surface is touched. ### Impacted Components - `projects/miopen/CHANGELOG.md` only. ### Potential Side Effects None expected. This does not change any runtime behavior; it only documents changes that already shipped on `release/therock-7.14`. ### Mitigation Steps None required beyond normal review — a reviewer familiar with the 7.14 release can spot-check a few entries against the linked commit hashes/PR numbers for accuracy. ## Test Plan Not applicable (no code changes). Verification was done by cross-referencing every added bullet against a specific commit hash and PR number pulled from `git log`, and by diffing the edited section against `git show remotes/origin/release/therock-7.14:projects/miopen/CHANGELOG.md` to confirm no existing entry was duplicated or altered. ## Test Result All 21 new bullets map 1:1 to a real commit/PR in the `therock-7.13..release/therock-7.14` range for `projects/miopen`. No pre-existing changelog text was modified — `git diff` shows a pure addition (24 insertions, 0 deletions). Two borderline commits (`a67834413e8` "Standalone build", `7c7cd5774a9` "enable warnings about deprecated API/features") were deliberately left out pending author judgment on whether they warrant an entry.
JIRA ID: ALMIOPEN-2312 ## Motivation `projects/miopen/CHANGELOG.md`'s "MIOpen 3.5.2 for ROCm 7.14.0" section only had 4 entries, but 68 commits touched `projects/miopen` between the `therock-7.13` tag and `release/therock-7.14`. Auditing that commit range (via `git log --oneline therock-7.13..remotes/origin/release/therock-7.14 -- projects/miopen`) found 18 additional user-facing changes — new gfx950/gfx1250 architecture support, SystemDB refreshes, AI-heuristic improvements, several correctness/crash bug fixes, and OpenCL solver removals — that were never documented. This PR backfills those entries so the 7.14.0 changelog accurately reflects what shipped. ## Technical Details Added 24 lines to the existing "MIOpen 3.5.2 for ROCm 7.14.0" section of `projects/miopen/CHANGELOG.md`, merging new bullets into the existing `Added`/`Changed`/`Resolved Issues` sub-sections (no new entries were duplicated or reworded) and adding a new `Removed` sub-section for solver removals. No existing content was altered. Each new bullet was traced to a specific commit + PR: | Entry | Commit | PR | |---|---|---| | Added gfx950 (MI350X/MI355X) SystemDB | `60d8e1e5175` | #8081 | | Added initial gfx1250 architecture support | `d3f057b2f22` | #7587 | | Added JSON-formatted performance logs (`MIOPEN_PERFORMANCE_LOGS`) | `ffbde871260` | #5736 | | Refreshed gfx942 SystemDB | `faf754a9ef6` | #8174 | | Refined 2D AI heuristics (Tunanet/KTN) for gfx942 & gfx950 | `c24ad415884` | #7109 | | Fixed thread-safety issue in AI-heuristic model caches | `cd957402309` | #9149 (cherry-pick of #9131) | | Fixed CK grouped-conv solvers not always registered in host | `c971d54d7ae` | #9105 | | Fixed legacy CK buffer resource descriptor (gfx1250) | `1dee011b8b0` | #8758 | | Removed stale gfx942 SystemDB naive-solver entries | `f42668f6316` | #7028 | | Fixed softmax NaN when `beta=0` | `fb8157c8512` | #8426 | | Fixed BN backward spatial variant 1 vectorized memory access bug | `1351a7e8b10` | #6765 | | Fixed use-after-free of `kernel_id` in `InitInvokerFactoryNCHW` | `394ebffb50d` | #7922 | | Fixed SoftMarginLoss forward reduce-kernel kernarg corruption | `1e65a31570f` | #8065, #8135 | | Fixed build with `MIOPEN_ENABLE_AI_KERNEL_TUNING=OFF` | `e6288c3ca61` | #6800 | | Fixed install RPATH for MIOpenDriver / CK backend libs | `ef8a4cf2f93` | #7485 | | Fixed dangling lambda causing log-level 6 segfault | `190f9791737` | #6953 | | Enabled Windows long-path support for `addkernels` | `1f9d69724cf` | #7387 | | Removed OCL solver `ConvOclDirectFwd1x1` | `126c720cadd` | #6491 | | Removed OCL solver `conv_ocl_dir2Dfwdgen` | `2d499ddf709` | #6633 | | Removed OCL solvers `ConvOclBwdWrW1x1`, `ConvOclBwdWrW2` | `b5d43a33dc6` | #6859 | | Removed OCL solvers `OpTensorLeadingOnes`, `Op5dTensorGeneric` | `8642c83b7de` | #7347 | The remaining ~40 commits in the 68-commit range were deliberately excluded (tests/CI/build/docs/internal refactors/1:1 OCL→HIP swaps with no behavior change, per each commit's own stated scope) and are not represented here. ## Risk Assessment **Risk Level: 🟢 Low** This is a documentation-only change to a single Markdown file. No code, build configuration, tests, or public API surface is touched. ### Impacted Components - `projects/miopen/CHANGELOG.md` only. ### Potential Side Effects None expected. This does not change any runtime behavior; it only documents changes that already shipped on `release/therock-7.14`. ### Mitigation Steps None required beyond normal review — a reviewer familiar with the 7.14 release can spot-check a few entries against the linked commit hashes/PR numbers for accuracy. ## Test Plan Not applicable (no code changes). Verification was done by cross-referencing every added bullet against a specific commit hash and PR number pulled from `git log`, and by diffing the edited section against `git show remotes/origin/release/therock-7.14:projects/miopen/CHANGELOG.md` to confirm no existing entry was duplicated or altered. ## Test Result All 21 new bullets map 1:1 to a real commit/PR in the `therock-7.13..release/therock-7.14` range for `projects/miopen`. No pre-existing changelog text was modified — `git diff` shows a pure addition (24 insertions, 0 deletions). Two borderline commits (`a67834413e8` "Standalone build", `7c7cd5774a9` "enable warnings about deprecated API/features") were deliberately left out pending author judgment on whether they warrant an entry.
JIRA ID: ROCM-27477
Motivation
GetFdeepModelandGetModelinprojects/miopen/src/conv/heuristics/ai_heuristics.cppcacheloaded AI heuristic models in
static std::mapcontainers with no synchronization. Whenconvolutions run concurrently from multiple threads on the same device (e.g. per-thread MIOpen
handles), two callers can race on the same cache key: both find the key missing, both build a
model, and both insert. In
GetFdeepModelthe second insert (models[key] = std::move(model))move-assigns over the first thread's
unique_ptr<fdeep::model>, destroying the model object whilethe first thread is still using the reference it was handed — a use-after-free that segfaults deep
inside
fdeep's prediction code.This was surfaced by the PyTorch test
test_cuda.py::TestCuda::test_cudnn_multiple_threads_same_deviceon ROCm 7.14.0-RC1 / gfx942 (MI300 series), with an observed crash rate of roughly 3-17% per
triggering run. The exposing commit (#7109, which enabled the 2D AI heuristic default path for
gfx942/gfx950) predates the last known-good nightly, but the underlying missing-lock defect itself
is older and not a regression introduced within the RC window. The race is warm-up-bounded: it can
only occur during first concurrent population of a given cache key, since entries are never
overwritten afterward.
A runtime workaround exists without rebuilding:
MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS_AI_HEUR=0disables the AI heuristic pathentirely, but that is not a substitute for fixing the underlying race.
Technical Details
The fix mirrors the already-correct sibling
GetCandidateSelectionModelinai_candidate_selection.cpp, which guards its structurally identical cache with astatic std::mutex+std::lock_guard.GetModel(arch, solver): added astatic std::mutex models_mutexand moved themodels.find(solver)lookup inside astd::lock_guardthat also covers the insert-on-misspath, so find+insert is atomic with respect to other callers. The function now returns the
cached
std::shared_ptr<Model>by value while still holding the lock, which is safe sinceshared_ptrcopies are cheap and the returned value keeps the model alive independently of themap.
GetFdeepModel(path, key): added astatic std::mutex models_mutexguarding find+insert.Critically, the insertion path was changed from
models[key] = std::move(model); return ref;(which risks a move-assign clobbering an existing entry on a race, and previously took the
returned reference before the map write) to
it = models.emplace(key, std::move(model)).first;followed by
return *(it->second);. Because entries are only ever inserted once and neverreassigned, a reference into the map obtained after the lock is released remains valid for the
lifetime of the process — the lock only needs to serialize the population itself, not every
subsequent read.
No public API, call signatures, or heuristic selection logic changed; this is purely a
thread-safety fix around existing caching behavior.
Risk Assessment
Risk Level: 🟢 Low
The change is a single-file, 27-line diff that adds standard mutex-guarded caching to two
functions, following the exact pattern already used correctly elsewhere in the same file. It does
not alter model selection, loading logic, or any public interface.
Impacted Components
projects/miopen/src/conv/heuristics/ai_heuristics.cpp:GetModel(used by the general AIheuristic model cache) and
candidate_selection::GetFdeepModel(used by the AI candidateselection path). Any convolution solver selection path that goes through the AI heuristics
(currently defaulted on for gfx942/gfx950 2D convolutions) is indirectly affected.
Potential Side Effects
GetModel/GetFdeepModel, not only on cache misses. Since these are in-process, uncontended locks afterwarm-up (steady-state calls hit the fast
findpath while holding a brief, uncontended lock),the performance impact is expected to be negligible relative to the cost of running the model
itself, but reviewers should watch for any contention under very high thread-count concurrent
convolution workloads.
concurrency safety of caching is affected.
Mitigation Steps
GetCandidateSelectionModel) already in productionin the same file, minimizing the chance of a novel defect.
MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS_AI_HEUR=0)remains available as an immediate mitigation if an unexpected issue surfaces post-merge, without
requiring a revert or rebuild.
possible if needed.
Test Plan
This will be validated manually on gfx942 hardware using the PyTorch repro test
(
test_cuda.py::TestCuda::test_cudnn_multiple_threads_same_device) that originally surfaced thecrash, running it repeatedly to build confidence that the segfault no longer reproduces.
In addition, unit tests have been added that reproduce the race condition. These have been run manually, and will also run as part of the standard MIOpen test category.
Test Result