tqcq/llama.cpp
0
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2025-07-19 17:17:40 +00:00
Code Issues Packages Projects Releases Wiki Activity

CUDA: fix race condition in MMQ stream-k fixup (#13299)

Browse Source
This commit is contained in:
Johannes Gäßler
2025-05-04 14:16:39 +02:00
committed by GitHub
parent 8afbd96818
commit 93c4e23905
1 changed files with 1 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
ggml/src/ggml-cuda/mmq.cuh
View File

@ -2958,6 +2958,7 @@ static __global__ void mul_mat_q_stream_k_fixup(
for (int j = threadIdx.y*WARP_SIZE + threadIdx.x; j < mmq_x; j += nwarps*WARP_SIZE) { for (int j = threadIdx.y*WARP_SIZE + threadIdx.x; j < mmq_x; j += nwarps*WARP_SIZE) {
ids_dst_shared[j] = ids_dst[col_low + j]; ids_dst_shared[j] = ids_dst[col_low + j];
} }
__syncthreads();
const int offset_dst = it*mmq_y; const int offset_dst = it*mmq_y;
dst += offset_dst; dst += offset_dst;