From 1b1dd65b836369e83cebef1fd21a5ec6808e3503 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Zolt=C3=A1n=20Lakatos?= <153429852+zsotakal@users.noreply.github.com> Date: Fri, 30 Jan 2026 08:22:54 +0100 Subject: [PATCH] fix undefined behaviour in softmax kernel (#3683) Co-authored-by: root [ROCm/composable_kernel commit: 565fea26455b8e4f78ac57ed64d6bd12e701a9c9] --- include/ck/tensor_operation/gpu/grid/gridwise_softmax.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_softmax.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_softmax.hpp index 96e13ac55c..a6fa04a824 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_softmax.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_softmax.hpp @@ -26,7 +26,7 @@ __global__ void kernel_softmax(const GridDesc_M_K in_grid_desc_m_k, AccDataType alpha, const InDataType* const __restrict__ p_in_value_global, AccDataType beta, - OutDataType* const __restrict__ p_out_value_global) + OutDataType* p_out_value_global) { GridwiseReduction::Run(in_grid_desc_m_k, out_grid_desc_m_k, @@ -91,7 +91,7 @@ struct GridwiseSoftmax_mk_to_mk AccDataType alpha, const InDataType* const __restrict__ p_in_value_global, AccDataType beta, - OutDataType* const __restrict__ p_out_value_global) + OutDataType* p_out_value_global) { if constexpr(SweepOnce) {