Skip to content

Commit 74503ef

Browse files
liangliangchangcursoragent
authored andcommitted
Fix Q6_K B-tile prefetch using wrong K offset in vec_dot
The pipelined load_ldmatrix for B0 omitted the k01 term, reading the wrong activation slice and corrupting Q6_K matmul results. Restores sanity-check accuracy on Qwen2.5 and Llama-2 while keeping the prefetch. Co-authored-by: Cursor <cursoragent@cursor.com>
1 parent 05f28d2 commit 74503ef

1 file changed

Lines changed: 1 addition & 1 deletion

File tree

ggml/src/ggml-cuda/mmq.cuh

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2867,7 +2867,7 @@ static __device__ __forceinline__ void vec_dot_q6_K_q8_1_mma(
28672867
constexpr int j_step = ntx*tile_C::J;
28682868

28692869
tile_B B0;
2870-
load_ldmatrix(B0, y_qs + 0, MMQ_TILE_Y_K);
2870+
load_ldmatrix(B0, y_qs + k01, MMQ_TILE_Y_K);
28712871
float dB0 = y_df[tile_C::get_j(0)*MMQ_TILE_Y_K + k01/QI8_1];
28722872

28732873
#pragma unroll

0 commit comments

Comments
 (0)