From 795e34d3d48be362bda5b0a0e787ee2bb67c4764 Mon Sep 17 00:00:00 2001 From: Jonathan Lichtner <195780826+JonathanLichtnerAMD@users.noreply.github.com> Date: Fri, 5 Jun 2026 21:04:58 +0000 Subject: [PATCH] [MIOpen] Fix SoftMarginLossForward reduce kernel kernarg corruption 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 --- .../solver/softmarginloss/forward_softmarginloss.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/miopen/src/solver/softmarginloss/forward_softmarginloss.cpp b/projects/miopen/src/solver/softmarginloss/forward_softmarginloss.cpp index 8c464f1b99d8..f86fa4ff2a1f 100644 --- a/projects/miopen/src/solver/softmarginloss/forward_softmarginloss.cpp +++ b/projects/miopen/src/solver/softmarginloss/forward_softmarginloss.cpp @@ -235,8 +235,11 @@ ConvSolution SoftMarginLossForward::GetSolution( size = (size + LOCAL_SIZE_REDUCE - 1) / LOCAL_SIZE_REDUCE; } + // The final ReduceSum kernel signature takes tensor_view_t<1>; passing the + // 5D o_tv would corrupt the kernarg buffer (80B written into 16B slot). + auto o_tv_1d = get_inner_expanded_tv<1>(deref(params.oDesc)); decltype(auto) kernel = handle_.Run(kernels[kernelCnt]); - kernel(reduce_in, params.o, size, o_tv); + kernel(reduce_in, params.o, size, o_tv_1d); if(profiling) { @@ -244,14 +247,11 @@ ConvSolution SoftMarginLossForward::GetSolution( (void)hipEventSynchronize(stop.get()); (void)hipEventElapsedTime(&elapsed, start.get(), stop.get()); - // Clean up - (void)hipEventDestroy(start.get()); - (void)hipEventDestroy(stop.get()); handle_.ResetKernelTime(); handle_.AccumKernelTime(elapsed); handle_.EnableProfiling(true); - }; + } }; }; }