Repository navigation
[MIOpen] Fix SoftMarginLossForward reduce kernel kernarg corruption (#8065) - #8135
Merged
JonathanLichtnerAMD merged 1 commit intoJun 8, 2026
Conversation
The final ReduceSum kernel signature takes tensor_view_t<1>, but the invoker was passing the 5D o_tv (80 bytes into a 16-byte slot), corrupting the kernarg buffer and producing deferred hipError 709 on Full/GPU_SoftMarginLossForward_FP32 MEAN/SUM cases (e.g. Test/82, dims:200x50 strides:1x200). Also drop the redundant manual hipEventDestroy calls in the profiling path; HipEventPtr's RAII deleter already handles destruction on scope exit, matching the canonical HipEventProfiler pattern. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
BrianHarrisonAMD
approved these changes
Jun 5, 2026
JonathanLichtnerAMD
deleted the
users/jlichtne/ROCMLIB-8065-softmarginloss-fwd-reduction-fix
branch
June 8, 2026 14:40
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Jul 15, 2026
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.
pschang-phy
pushed a commit
that referenced
this pull request
Jul 17, 2026
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.
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.
Motivation
Fixes #8065.
Full/GPU_SoftMarginLossForward_*withreduction_modeMEAN (2) or SUM (1)failed in the
-L fullCI sweep ongfx94X-dcgpu(MI300X) withhipError 709("context is destroyed"), blocking the rocm-systems → TheRock bump PRs.
Root cause: the reduction invoker in
projects/miopen/src/solver/softmarginloss/forward_softmarginloss.cpppassed atensor_view_t<5>(80 B) as the last argument to the finalReduceSumkernel,whose signature takes
tensor_view_t<1>(16 B). The 64 B overflow corruptedthe kernarg buffer; the kernel then used garbage stride/offset fields for
address computation, and the resulting wild GPU stores corrupted HSA pool
tracking. The deferred
hipError 709was raised later by theHIPErrorHandlergtest listener at test teardown.
Risk Assessment
Low. Five-line change confined to a single MIOpen solver
(
SoftMarginLossForward). The change replaces an incorrectly-sizedtensor_view_targument with the correctly-sized one — matching the kernel'sdeclared signature — and removes manual
hipEventDestroycalls that duplicatethe work
HipEventPtr's RAII deleter already performs (matching the canonicalHipEventProfilerpattern insrc/hipoc/hipoc_kernel.cpp). No public API,ABI, header, schema, or build-system changes. No algorithmic changes. The
reduction_mode = NONEpath is untouched.Technical Details
tensor_view_t<1>from the output descriptor and pass it to thefinal
ReduceSumkernel instead of the 5Do_tvused by the per-elementphase.
(void)hipEventDestroy(...)calls in theif(profiling)cleanup block —
HipEventPtr(aMIOPEN_MANAGE_PTR(hipEvent_t, hipEventDestroy)alias) destroys both events on scope exit.Test Plan
Manually reproduced and verified on
gfx942(MI300X) under AddressSanitizer(
-fsanitize=address -shared-libasan,HSA_XNACK=1) using the MIOpen gtestbinary. CI should handle the rest.
Test Result
Full/GPU_SoftMarginLossForward_FP32.*tests passed.CI will exercise the FP16 / BFP16 variants automatically.
Submission Checklist