From 26113d1dd388acd3ddc84c1af1f206116db3bd62 Mon Sep 17 00:00:00 2001 From: Yap Sok Ann Date: Mon, 24 Aug 2026 23:48:12 +0700 Subject: [PATCH] ggml-cuda: bind cublas handle to the backend stream in DSA attention (#2347) The DSA attention kernel used the shared cublas handle without binding it to the backend's stream, so its Q.K / P.V GEMMs ran on a different stream than the gather and softmax kernels. The softmax could then read the score buffer before the GEMM wrote it, picking up stale (NaN) values. --- ggml/src/ggml-cuda/dsa_attn.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/ggml/src/ggml-cuda/dsa_attn.cu b/ggml/src/ggml-cuda/dsa_attn.cu index dbdc5bca..57ad2ac1 100644 --- a/ggml/src/ggml-cuda/dsa_attn.cu +++ b/ggml/src/ggml-cuda/dsa_attn.cu @@ -344,6 +344,7 @@ bool ggml_cuda_dsa_attn_ext(ggml_backend_cuda_context & ctx, ggml_tensor * dst) (const float *)((const char *)Q->data + first*Q->nb[1]), q16.get()); } + CUBLAS_CHECK(cublasSetStream(ctx.cublas_handle(), ctx.stream())); CUBLAS_CHECK(cublasHgemmStridedBatched(ctx.cublas_handle(), CUBLAS_OP_T, CUBLAS_OP_N, indexer->ne[0], Q->ne[2], Q->ne[0], &alpha, k16.get(), K->ne[0], K->ne[0]*indexer->ne[0],