From f535d8d421c079ebeffd924e7154685fe1eaed58 Mon Sep 17 00:00:00 2001 From: Kawrakow Date: Fri, 17 Jul 2026 15:05:15 +0000 Subject: [PATCH] Fix race in indexer topk on CUDA --- ggml/src/ggml-cuda/indexer_topk.cu | 2 ++ 1 file changed, 2 insertions(+) diff --git a/ggml/src/ggml-cuda/indexer_topk.cu b/ggml/src/ggml-cuda/indexer_topk.cu index 14b12386b..a110a7ace 100644 --- a/ggml/src/ggml-cuda/indexer_topk.cu +++ b/ggml/src/ggml-cuda/indexer_topk.cu @@ -240,7 +240,9 @@ static __global__ void k_indexer_mask(int ne0, int ne1, int ne2, int ntopk, int if (i1 < ne11) { for (int j = threadIdx.x; j < ne0; j += blockDim.x) d[j] = inf; + __syncthreads(); for (int j = threadIdx.x; j < ntopk; j += blockDim.x) d[i[j]] = zero; + __syncthreads(); for (int j = threadIdx.x; j < ne0; j += blockDim.x) d[j] += m[j]; } else { for (int j = threadIdx.x; j < ne0; j += blockDim.x) d[j] = m[j];