From: Aman Gupta <redacted>
Date: Wed, 22 Oct 2025 04:33:08 +0000 (+0800)
Subject: CUDA: fix bug in topk-moe softmax (llama/16711)
X-Git-Tag: upstream/0.9.4.185~94
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=db51a43467eb7e0b410ce06265e9608eb16a267c;p=pkg%2Fggml%2Fsources%2Fggml

CUDA: fix bug in topk-moe softmax (llama/16711)
---

diff --git a/src/ggml-cuda/topk-moe.cu b/src/ggml-cuda/topk-moe.cu
index d782ad94..e28c810a 100644
--- a/src/ggml-cuda/topk-moe.cu
+++ b/src/ggml-cuda/topk-moe.cu
@@ -141,7 +141,7 @@ __launch_bounds__(4 * WARP_SIZE, 1) __global__ void topk_moe_cuda(const float *
         wt_sum              = warp_reduce_sum(wt_sum);
         const float inv_sum = 1.0f / wt_sum;
 
-        for (int i = threadIdx.x; i < n_expert_used; i += WARP_SIZE) {
+        for (int i = 0; i < experts_per_thread; i++) {
             output_weights[i] *= inv_sum;
         }
     }