Repository navigation
[MIOpen] Fix dangling lambda capture causing log-level 6 segfault - #6953
Merged
Merged
Conversation
amd-ahyttine
marked this pull request as ready for review
May 6, 2026 06:09
vpietila-amd
approved these changes
May 6, 2026
vpietila-amd
left a comment
Contributor
There was a problem hiding this comment.
Looks good to me. Nice catch.
1 task done
aledudek
pushed a commit
that referenced
this pull request
May 20, 2026
) ## Motivation This PR fixes "Segmentation fault (core dumped)" when running MIOpenDriver. Example: "MIOPEN_FIND_MODE=2 MIOPEN_FIND_ENFORCE=1 MIOPEN_LOG_LEVEL=6 ./MIOpenDriver convfp16 --spatial_dim 2 --in_channels 24 --out_channels 12 --batchsize 4 --in_layout NHWC --out_layout NHWC --fil_layout NHWC --group_count 1 --forw 4 --pad_h 0 --pad_w 0 --in_h 112 --in_w 112 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --mode conv --search 1 --time 1". However when MIOPEN_LOG_LEVEL=5 it runs fine, no segmentation fault. The problem is limited to NHWC layout, direction backwards weights and MIOPEN_LOG_LEVEL>=6. ## Technical Details The level-6 run goes much deeper into Info2 logging and appears to crash during candidate-selection debug logging. Gdb confirms a host-side crash in memcpy while writing to an output stream, inside libMIOpenCKGroupedConv_gfx950.so during backward-weights invocation. That strongly points to a bad string/logging path in the CK grouped-conv runtime. The root bug: one invoker lambda captures kernel_id by reference to a temporary outer lambda capture, so at log level 6 (Info2) the kernel_name log dereferences a dangling string and crashes in memcpy. I’m applying a one-line fix to capture that string by value. const std::vector<Kernel>&) mutable { return [&kernel_id2 = kernel_id_, return [kernel_id2 = kernel_id_, split_k2 = split_k_, Crash cause is a dangling reference in a log-only path. In src/ck_impl/implicitgemm_ck_util.hpp, the NHWC WrW invoker lambda captured kernel_id like this: &kernel_id2 = kernel_id_. That reference points to an outer lambda capture that no longer exists when the invoker runs later. At MIOPEN_LOG_LEVEL=6, MIOPEN_LOG_I2("kernel_name = " << kernel_id2) tries to print that dangling string, causing a host segfault (memcpy in std::ostream), which matches your gdb backtrace. At MIOPEN_LOG_LEVEL=5, that Info2 log is not executed, so the bug stays hidden. I fixed it by capturing kernel_id2 by value instead of by reference: Updated src/ck_impl/implicitgemm_ck_util.hpp: return [&kernel_id2 = kernel_id_, ...] to return [kernel_id2 = kernel_id_, ...] Note that other occurrences of kernel_id2 in the file have the correct "kernel_id2= kernel_id_" version, instead of this faulty "&kernel_id2 = kernel_id_" version. ## Test Plan As validation test the command mentioned in the motivation (by AI). Also rebuild MIOPenDriver and running 1000 command performance comparison that before the fix resulted in several segfaults and other crashes on MIOPEN_LOG_LEVEL=6. ## Test Result Validation done: Rebuilt MIOpenDriver and libMIOpenCKGroupedConv_gfx950.so. Re-ran your exact reproducer with MIOPEN_LOG_LEVEL=6. Result: no segfault, run completes and verifies OK. Rebuilt MIOPenDriver manually, all 1000 MIOpenDriver commands run now succesfully. ## Submission Checklist - [v] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: root <root@smci355-ccs-aus-m02-05.cs-aus.dcgpu>
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
This PR fixes "Segmentation fault (core dumped)" when running MIOpenDriver. Example:
"MIOPEN_FIND_MODE=2 MIOPEN_FIND_ENFORCE=1 MIOPEN_LOG_LEVEL=6 ./MIOpenDriver convfp16 --spatial_dim 2 --in_channels 24 --out_channels 12 --batchsize 4 --in_layout NHWC --out_layout NHWC --fil_layout NHWC --group_count 1 --forw 4 --pad_h 0 --pad_w 0 --in_h 112 --in_w 112 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --mode conv --search 1 --time 1".
However when MIOPEN_LOG_LEVEL=5 it runs fine, no segmentation fault.
The problem is limited to NHWC layout, direction backwards weights and MIOPEN_LOG_LEVEL>=6.
Technical Details
The level-6 run goes much deeper into Info2 logging and appears to crash during candidate-selection debug logging.
Gdb confirms a host-side crash in memcpy while writing to an output stream, inside libMIOpenCKGroupedConv_gfx950.so during backward-weights invocation. That strongly points to a bad string/logging path in the CK grouped-conv runtime.
The root bug: one invoker lambda captures kernel_id by reference to a temporary outer lambda capture, so at log level 6 (Info2) the kernel_name log dereferences a dangling string and crashes in memcpy. I’m applying a one-line fix to capture that string by value.
Crash cause is a dangling reference in a log-only path.
In src/ck_impl/implicitgemm_ck_util.hpp, the NHWC WrW invoker lambda captured kernel_id like this: &kernel_id2 = kernel_id_.
That reference points to an outer lambda capture that no longer exists when the invoker runs later.
At MIOPEN_LOG_LEVEL=6, MIOPEN_LOG_I2("kernel_name = " << kernel_id2) tries to print that dangling string, causing a host segfault (memcpy in std::ostream), which matches your gdb backtrace.
At MIOPEN_LOG_LEVEL=5, that Info2 log is not executed, so the bug stays hidden.
I fixed it by capturing kernel_id2 by value instead of by reference:
Updated src/ck_impl/implicitgemm_ck_util.hpp:
return [&kernel_id2 = kernel_id_, ...]
to return [kernel_id2 = kernel_id_, ...]
Note that other occurrences of kernel_id2 in the file have the correct "kernel_id2= kernel_id_" version, instead of this faulty "&kernel_id2 = kernel_id_" version.
Test Plan
As validation test the command mentioned in the motivation (by AI).
Also rebuild MIOPenDriver and running 1000 command performance comparison that before the fix resulted in several segfaults and other crashes on MIOPEN_LOG_LEVEL=6.
Test Result
Validation done:
Rebuilt MIOpenDriver and libMIOpenCKGroupedConv_gfx950.so.
Re-ran your exact reproducer with MIOPEN_LOG_LEVEL=6.
Result: no segfault, run completes and verifies OK.
Rebuilt MIOPenDriver manually, all 1000 MIOpenDriver commands run now succesfully.
Submission Checklist