[jdcampbe]

Running any NCHW grouped CK xdlops convolution with kernel-name logging enabled (e.g. MIOPEN_PERFORMANCE_LOGS=4) deterministically SIGSEGVs inside
MIOpen's invoker. This was originally surfaced as a CK / large-tensor regression in a perf-eval run, but root-causing showed the fault is host-side in
MIOpen and reproduces independent of data type or tensor size — any NCHW grouped CK conv that runs AddKernelToJsonAccumulator will crash.

This PR fixes the use-after-free so the NCHW invoker matches the by-value capture used by the three sibling invoker factories in the same file.

Technical Details

In projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp, InitInvokerFactoryNCHW captured its const std::string& kernel_id parameter by
pointer:

result.invoker_factory = [kernel_id_ = &kernel_id, ...]  // dangling once function returns
    return [kernel_id2 = kernel_id_, ...]                // kernel_id2 is const std::string*
        ...
        MIOPEN_LOG_I2("kernel_name = " << kernel_id2);            // prints stale stack address
        AddKernelToJsonAccumulator(*kernel_id2, elapsed, false);  // dereferences dangling ptr -> SIGSEGV

The invoker_factory (and its nested invoker lambda) are stored in the returned ConvSolution and run long after InitInvokerFactoryNCHW has
returned, by which time the referent of the kernel_id parameter has been destroyed. The deref is gated on IsProfilingEnabled() && IsLoggingKernel(), so it fires whenever kernel-name JSON logging is on during execution.

The other three invoker factories in this file (NHWC fwd, NHWC bwd/wrw) already capture kernel_id by value (lines 690, 1138, 1210). An earlier
dereference-error fix corrected the *kernel_id2 deref in some sibling paths but missed this NCHW one.

This 3-line change:

• Capture kernel_id by value: kernel_id_ = kernel_id (line 981).
• Move the owned string into the inner lambda: kernel_id2 = std::move(kernel_id_) (line 991).
• Drop the dereference: AddKernelToJsonAccumulator(kernel_id2, elapsed, false) (line 1079).

As a beneficial side-effect, MIOPEN_LOG_I2("kernel_name = " << kernel_id2) now prints the actual kernel name instead of the stale pointer address.

The remaining &kernel_id capture in FindConvPtrByID (line 165) is a synchronous, non-escaping std::find_if lambda and is safe.

Test Plan

• Re-run the original reproducer (e.g. MIOpenDriver conv -n 1 -c 512 -H 912 -W 1144 -k 512 -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1 -iter 10) with MIOPEN_FIND_MODE=1 MIOPEN_FIND_ENFORCE=4 MIOPEN_LOG_LEVEL=7 MIOPEN_PERFORMANCE_LOGS=4 on gfx950 and confirm it no longer
  SIGSEGVs.
• Re-run the broader perf-eval shape-coverage shard that produced ~69 segfaults on this same code path and confirm those crashes are gone.
• Spot-check the MIOPEN_LOG_I2("kernel_name = …") log line now prints the kernel name string (e.g.
  DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<...>) rather than a 0x7fff… stack address.
• Existing MIOpen gtests covering NCHW grouped CK conv paths should continue to pass.

Test Result

• Static analysis: confirmed no remaining *kernel_id2 dereferences and no escaping &kernel_id captures in implicitgemm_ck_util.hpp; all four
  invoker factories now capture kernel_id by value.
• Validated at runtime through a battery of tests.

Risk Assessment

Risk level: Low

• Impacted components: projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp, specifically the NCHW invoker-factory path used by the
  ConvHipImplicitGemmGroup{Fwd,Bwd,Wrw}Xdlops and ConvHipImplicitGemm3DGroup{Fwd,Bwd,Wrw}Xdlops solvers when the problem is in default (NCHW/NCDHW)
  layout. No CK code, no kernel code, no public API changes. Other MIOpen layouts and solvers are unaffected.
• Potential side effects:
  • Functional: removes a guaranteed use-after-free; no behavior change on the happy path besides the log line now printing the kernel name instead of
    a stale address.
  • Performance: one extra std::string copy per InitInvokerFactoryNCHW call (a one-time setup cost during solution build), then a std::move into
    the inner lambda — negligible vs the kernel launch and CK arg setup cost. Per-invocation cost on the hot path is unchanged.
  • ABI/API: none.
• Mitigation steps:
  • Change is minimal (3 lines) and matches the established by-value pattern already in use by the three sibling invoker factories in the same file, so
    review surface is small.
  • The defect is straightforwardly reproducible with MIOPEN_PERFORMANCE_LOGS on any NCHW grouped CK conv, making the fix easy to verify pre-merge.
  • If a regression surfaces, the change is trivially revertible.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[BrianHarrisonAMD]

LGTM.

Minor comments, but I don't think it's blocking.

[JonathanLichtnerAMD]

LGTM. Would it be worth adding a test for this, i.e., it could either reproduce the crash or ASAN could detect the issue? I'll leave that as optional but having a test means that ASAN and other sanitizers will validate the code path against future changes.

[yassinsolim]

LGTM. The lifetime fix is correct: kernel_id is now owned by the factory/invoker instead of captured by pointer to a local. Brian’s move-consistency comments seem reasonable but non-blocking to me.