From efc74f8bceee40699bafdcaee2042fe00d4c852a Mon Sep 17 00:00:00 2001 From: jdcampbe Date: Sun, 31 May 2026 00:57:44 -0500 Subject: [PATCH] [MIOpen] Fix use-after-free of kernel_id in InitInvokerFactoryNCHW InitInvokerFactoryNCHW captured the kernel_id string parameter by pointer (kernel_id_ = &kernel_id) instead of by value. The invoker_factory and its nested invoker lambda are stored in the returned ConvSolution and executed long after InitInvokerFactoryNCHW returns, by which point the referent of the `const std::string& kernel_id` parameter has been destroyed. At invoke time the dangling pointer was dereferenced via AddKernelToJsonAccumulator(*kernel_id2, ...) (gated on profiling + IsLoggingKernel), reading a freed std::string and crashing with SIGSEGV. The accompanying MIOPEN_LOG_I2("kernel_name = " << kernel_id2) also printed the stale pointer address instead of the kernel name. All three sibling invoker factories in this file already capture kernel_id by value; only the NCHW forward path used a pointer (and the earlier dereference-error fix corrected the other paths but missed this one). Capture kernel_id by value here too so the lambda owns the string for its full lifetime, drop the now-unnecessary dereference, and let the log print the actual kernel name. This is reached whenever an NCHW grouped CK xdlops convolution is executed with kernel logging enabled (e.g. MIOPEN_PERFORMANCE_LOGS), independent of data type or tensor size. Co-Authored-By: Claude Opus 4.8 --- projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp b/projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp index 8775d9edf3d5..6d77e06981e9 100644 --- a/projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp +++ b/projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp @@ -933,7 +933,7 @@ ConvSolution InitInvokerFactoryNCHW(const ExecutionContext& ctx, internal::MakeTaggedTransposeInstances( result, ctx, problem, ck_args, input1_op, input2_op, output_op, _ck_buff_des); - result.invoker_factory = [kernel_id_ = &kernel_id, + result.invoker_factory = [kernel_id_ = kernel_id, split_k_ = split_k, ck_args_ = std::move(ck_args), sh_conv_ptr_ = std::shared_ptr{std::move(*ptr_iter)}, @@ -943,7 +943,7 @@ ConvSolution InitInvokerFactoryNCHW(const ExecutionContext& ctx, output_init_tr_inst_ = std::move(_output_init_tr_inst), ck_buff_des_ = _ck_buff_des](const std::vector& kernels) mutable { - return [kernel_id2 = kernel_id_, + return [kernel_id2 = std::move(kernel_id_), split_k2 = split_k_, kernels, ck_args2 = std::move(ck_args_), @@ -1031,7 +1031,7 @@ ConvSolution InitInvokerFactoryNCHW(const ExecutionContext& ctx, // Kernel logging for CK kernels if(IsLoggingKernel()) { - AddKernelToJsonAccumulator(*kernel_id2, elapsed, false); + AddKernelToJsonAccumulator(kernel_id2, elapsed, false); } handle.ResetKernelTime(); handle.AccumKernelTime(elapsed);