Repository navigation
fix(miopen): validate work dimensions on solution deserialization - #12242
Merged
JonathanLichtnerAMD merged 3 commits intoSep 21, 2026
Merged
Conversation
Attacker-controlled local_work_dims/global_work_dims are deserialized as unbounded std::vector<size_t> from a serialized Solution blob (reachable via the public miopenLoadSolution API) and later copied into a fixed std::array<size_t,3> in the HIPOCKernel constructor, guarded only by assert() -- which compiles out under NDEBUG, leaving a heap out-of-bounds write in release builds. Harden the path at two layers: - hipoc_kernel.hpp: replace the assert() guards in the HIPOCKernel ctor with real MIOPEN_THROW(miopenStatusBadParm) checks, bounded against the destination array capacity (ldims.size()/gdims.size()) so the std::copy can never overrun. - solution.cpp/solution.hpp: validate work-dims size in SerializedSolutionKernelInfo::from_json via a new ValidateSerializedWorkDims helper, rejecting a malicious blob at load time (miopenStatusInvalidValue) before any kernel is constructed. - test/gtest/solution_load_work_dims.cpp: GPU-free regression test asserting oversized (>3) and empty work-dims are rejected and 1..3 are accepted. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
iq136boy
approved these changes
Sep 17, 2026
averinevg
reviewed
Sep 17, 2026
comgr.cpp uses assert() but only got <cassert> transitively via miopen/kernel.hpp -> hipoc_kernel.hpp. The previous commit removed that transitive include from hipoc_kernel.hpp (it was the only local user of assert() there), which broke comgr.cpp's build. Include <cassert> directly instead of relying on an implicit transitive dependency.
…OCKernel By the time HIPOCKernel's constructor is reached, work-dimension vectors should already have been validated at the relevant trust boundary (e.g. Solution deserialization). A vector that fails this check here is an internal invariant violation rather than a bad caller-supplied parameter, so use miopenStatusInternalError instead of miopenStatusBadParm. Addresses review feedback on PR #12242.
JonathanLichtnerAMD
deleted the
users/jlichtne/rocm-31082-workdims-oob
branch
September 21, 2026 17:35
Contributor
TheRock Submodule Bump ActivityNewest first
|
evedovelli
pushed a commit
that referenced
this pull request
Sep 23, 2026
…2242) JIRA ID: ROCM-31082 ## Motivation `miopenLoadSolution` deserializes kernel work-dimension vectors (`local_work_dims` / `global_work_dims`) from a serialized `Solution` blob and copies them into a fixed `std::array<size_t, 3>` inside the `HIPOCKernel` constructor. The only bounds check on that copy was a debug-only `assert`, which is compiled out in release builds. A blob with more than 3 entries in either array is attacker/corruption-controlled input that overruns the destination array — an out-of-bounds write. This PR closes that gap. ## Technical Details - `miopen/src/solution.cpp`: added `ValidateSerializedWorkDims()`, called from `SerializedSolutionKernelInfo::from_json` right after the work-dims vectors are parsed, before any kernel object is constructed. Throws `miopenStatusInvalidValue` if the vector is empty or has more than 3 entries. - `miopen/src/include/miopen/hipoc_kernel.hpp`: replaced the release-mode-inert `assert(!local_dims.empty() && local_dims.size() <= 3)` (and the `global_dims` equivalent) in the `HIPOCKernel` constructor with `MIOPEN_THROW(miopenStatusInternalError, ...)` checks against the actual array capacity (`ldims.size()` / `gdims.size()`), so the same invariant is enforced in release builds as a second line of defense for any other caller path into this constructor. By this point the work dims should already be validated at the deserialization boundary above, so a violation here is an internal invariant break rather than a bad caller-supplied parameter. - `miopen/src/include/miopen/solution.hpp`: declared and documented the new `ValidateSerializedWorkDims` helper (`MIOPEN_INTERNALS_EXPORT`) so it's reusable and independently testable. - Added `miopen/test/gtest/solution_load_work_dims.cpp`, a CPU-only regression test exercising the helper directly with an oversized vector, an empty vector, and valid 1-3 element vectors — no GPU or full solution round-trip needed. - `miopen/src/comgr.cpp`: added a direct `#include <cassert>`. This file uses `assert()` but was only getting `<cassert>` transitively via `miopen/kernel.hpp` -> `hipoc_kernel.hpp`; removing that transitive include above broke its build (caught by CI), so it now includes `<cassert>` directly instead of relying on an implicit dependency. ## Risk Assessment The change only adds input validation (bounds checks that throw on invalid input) on a narrow deserialization path and a kernel constructor, plus an unrelated one-line include fix; it does not alter behavior for any valid work-dimensions input, so blast radius is small. **Risk Level: 🟢 Low** ### Impacted Components - `miopen/src/solution.cpp` (`SerializedSolutionKernelInfo::from_json`) - `miopen/src/include/miopen/solution.hpp` - `miopen/src/include/miopen/hipoc_kernel.hpp` (`HIPOCKernel` constructor) - `miopen/src/comgr.cpp` (missing direct `<cassert>` include) - New test: `miopen/test/gtest/solution_load_work_dims.cpp` ### Potential Side Effects None expected for well-formed solutions (1-3 work dimensions, the only valid case today). Any caller that was previously relying on the debug-only assert being absent in release builds to pass out-of-range dimensions would now get a thrown exception instead of undefined behavior — this is the intended fix, not a regression. ### Mitigation Steps Added a dedicated regression test covering the empty, oversized, and valid input cases. No feature flag or staged rollout needed given the contained, input-validation-only nature of the change. ## Test Plan Added `solution_load_work_dims.cpp`, a new CPU-only GTest target that calls `miopen::ValidateSerializedWorkDims` directly with an oversized vector (4 entries), an empty vector, and valid 1/2/3-entry vectors, asserting throw vs. no-throw and checking the thrown status code. ## Test Result CI ran on this PR and caught the `comgr.cpp` transitive-include break described above; that has been fixed in a follow-up commit. CI is expected to build and execute `solution_load_work_dims.cpp` as part of the standard gtest suite on the next run. --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
mattwill-amd
pushed a commit
that referenced
this pull request
Oct 14, 2026
JIRA ID: ALMIOPEN-2913 ## Motivation Only 5 of the 30 commits that landed against `projects/miopen` since the 3.6.2 version bump (`2d20a8f775e`, #12211) were reflected in `CHANGELOG.md`. The remaining 25 were reviewed and the 15 externally-visible ones documented, so the ROCm 10.2.0 section accurately reflects what shipped. 14 entries added (two of the fixes share one entry), plus a duplicate-heading fix. ### Added * [Conv] NHWC (channel-last) support in the GEMM convolution solvers, with new im2col/col2im kernels (#8693) * [Conv] Cross-architecture AI heuristics for immediate mode (`MIOPEN_FIND_MODE=2`), extending data-driven selection to Navi (#9680) * [Conv] Missing 3x3 stride-1 pad-1 tile-28 CK depthwise backward-data instance — ALMIOPEN-2423 (#11271) ### Changed * [Conv] `ConvHipConv` held out of the AI perf-config picker until its model is retrained — ALMIOPEN-2494 (#13337) * [Conv] CK byte-span gate lifted for forward and backward-data, now that CK fixes the 32-bit offset overflow itself — ALMIOPEN-2645 (#12396) ### Optimized * [Conv] `ConvDirectNaiveConv`: per-direction kernel files, float accumulators, spatial tiling, bfloat16 path — ALMIOPEN-2485/2486 (#7529) * Kernel cache reuses cached SQLite connections instead of reopening per load/store — ALMIOPEN-2598 (#11933) ### Resolved Issues * [Conv] Incorrect backward-weights results from the gfx950 direct kernel at `waves_q == 1` (missing cross-lane ordering on LDS staging reuse) — ROCM-31508 (#12406) * [Conv] Up to 8x slower hipconv `direct`/`direct_l1`/`direct_wgrad` under ROCm 10.1's full AMDGCN LTO default; fixed with `-fno-offload-lto` — ROCM-31334 (#12328) * [Conv] gfx942 multi-node training throughput regression from backward GEMM on small-batch 1x1 NHWC shapes — ROCM-31155 (#13232) * [Conv] Invalid CK perf-config entries in the gfx1201 and gfx1100 system databases — ROCM-29026 (#12088, #12117) * [BatchNorm] JIT compilation failures on non-GFX9 targets; Welford DPP assembly restricted to GFX9 — ALMIOPEN-2838 (#12889) * Out-of-bounds write in `miopenLoadSolution` when a serialized solution carries more than three work dimensions — ROCM-31082 (#12242) * `uint64_t` typedef conflict between `miopen_cstdint.hpp` and system `<stdint.h>` on LP64 Linux (#12623) ## Technical Details Single-file change to `projects/miopen/CHANGELOG.md`. The 10.2.0 section had **two** `### Resolved Issues` headings. The two entries under the trailing one (`ConvHipConv` gfx1250 crash/LDS race, and the perf-db reload config loss) were merged into the first, and the trailing heading removed — this accounts for the 2 removed bullets in the diff. Section order is now `Resolved Issues → Added → Changed → Optimized → Removed`, matching the 3.6.1 section. `Optimized` is new to this release section. Already documented, so untouched: #12874 (reduce int32 rejection), #12000 / #13016 / #13213 (hipconv v0.3.1 / v0.3.2 / v0.4.0), #11761 (gfx803 removal). ## Notes Deliberately excluded as not externally visible (10 commits): * **hipDNN convolution forwarding (#12276)** — `MIOPEN_ENABLE_HIPDNN_WRAPPER` defaults to `OFF` (`projects/miopen/CMakeLists.txt:162`), so no shipped build compiles the forwarding code and the runtime `MIOPEN_HIPDNN_FORWARDING` flag has nothing to enable. * **Test-only**: #13289 (gfx1250 depthwise test applicability, ALMIOPEN-2841), #11610 (hipDNN parity harness), #12660 (Jenkins status post). * **Internal kernel-math cleanup, generated ISA unchanged**: #13312 (drop direct `__ocml` calls), #13190 (remove broken single-argument `__half` overloads). * **Add/revert pairs, net no change**: #12466 + #13066 (gfx1101 Navi system DBs), #12496 + #12800 (gfx1250 in the dev image arch lists). ## Risk Assessment Documentation-only change to a single Markdown file; no code, build, or runtime impact. **Risk Level: 🟢 Low** ## Test Plan None required — Markdown-only. Verified that the entry topics are not already documented elsewhere in the file, that no content was lost when merging the duplicate heading, and that every new entry maps to a commit in `2d20a8f775e..origin/develop -- projects/miopen`. ## Test Result 16 bullets added, 2 relocated — 14 net-new entries covering 15 commits. Duplicate heading removed; section ordering matches 3.6.1. Co-authored-by: Claude <noreply@anthropic.com>
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
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.
JIRA ID: ROCM-31082
Motivation
miopenLoadSolutiondeserializes kernel work-dimension vectors(
local_work_dims/global_work_dims) from a serializedSolutionblob andcopies them into a fixed
std::array<size_t, 3>inside theHIPOCKernelconstructor. The only bounds check on that copy was a debug-only
assert,which is compiled out in release builds. A blob with more than 3 entries in
either array is attacker/corruption-controlled input that overruns the
destination array — an out-of-bounds write. This PR closes that gap.
Technical Details
miopen/src/solution.cpp: addedValidateSerializedWorkDims(), called fromSerializedSolutionKernelInfo::from_jsonright after the work-dims vectorsare parsed, before any kernel object is constructed. Throws
miopenStatusInvalidValueif the vector is empty or has more than 3entries.
miopen/src/include/miopen/hipoc_kernel.hpp: replaced the release-mode-inertassert(!local_dims.empty() && local_dims.size() <= 3)(and theglobal_dimsequivalent) in the
HIPOCKernelconstructor withMIOPEN_THROW(miopenStatusInternalError, ...)checks against the actual array capacity (
ldims.size()/gdims.size()),so the same invariant is enforced in release builds as a second line of
defense for any other caller path into this constructor. By this point the
work dims should already be validated at the deserialization boundary above,
so a violation here is an internal invariant break rather than a bad
caller-supplied parameter.
miopen/src/include/miopen/solution.hpp: declared and documented the newValidateSerializedWorkDimshelper (MIOPEN_INTERNALS_EXPORT) so it'sreusable and independently testable.
miopen/test/gtest/solution_load_work_dims.cpp, a CPU-only regressiontest exercising the helper directly with an oversized vector, an empty
vector, and valid 1-3 element vectors — no GPU or full solution round-trip
needed.
miopen/src/comgr.cpp: added a direct#include <cassert>. This file usesassert()but was only getting<cassert>transitively viamiopen/kernel.hpp->hipoc_kernel.hpp; removing that transitive includeabove broke its build (caught by CI), so it now includes
<cassert>directly instead of relying on an implicit dependency.
Risk Assessment
The change only adds input validation (bounds checks that throw on invalid
input) on a narrow deserialization path and a kernel constructor, plus an
unrelated one-line include fix; it does not alter behavior for any valid
work-dimensions input, so blast radius is small.
Risk Level: 🟢 Low
Impacted Components
miopen/src/solution.cpp(SerializedSolutionKernelInfo::from_json)miopen/src/include/miopen/solution.hppmiopen/src/include/miopen/hipoc_kernel.hpp(HIPOCKernelconstructor)miopen/src/comgr.cpp(missing direct<cassert>include)miopen/test/gtest/solution_load_work_dims.cppPotential Side Effects
None expected for well-formed solutions (1-3 work dimensions, the only valid
case today). Any caller that was previously relying on the debug-only assert
being absent in release builds to pass out-of-range dimensions would now get
a thrown exception instead of undefined behavior — this is the intended
fix, not a regression.
Mitigation Steps
Added a dedicated regression test covering the empty, oversized, and valid
input cases. No feature flag or staged rollout needed given the contained,
input-validation-only nature of the change.
Test Plan
Added
solution_load_work_dims.cpp, a new CPU-only GTest target that callsmiopen::ValidateSerializedWorkDimsdirectly with an oversized vector (4entries), an empty vector, and valid 1/2/3-entry vectors, asserting throw vs.
no-throw and checking the thrown status code.
Test Result
CI ran on this PR and caught the
comgr.cpptransitive-include breakdescribed above; that has been fixed in a follow-up commit. CI is expected to
build and execute
solution_load_work_dims.cppas part of the standard gtestsuite on the next run.