From 5b263dd83a5f906eddd10bc044051d7571097043 Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Sat, 3 Feb 2024 16:12:20 +0200
Subject: [PATCH] cuda : unroll Q*K^T loop

---
 ggml-cuda.cu | 1 +
 1 file changed, 1 insertion(+)

diff --git a/ggml-cuda.cu b/ggml-cuda.cu
index 67541a61e..dbd482239 100644
--- a/ggml-cuda.cu
+++ b/ggml-cuda.cu
@@ -6571,6 +6571,7 @@ static __global__ void flash_attn_ext_f16(
 
             // Q*K^T
             {
+#pragma unroll
                 for (int cc = 0; cc < C/16; ++cc) {
                     half16x16_acc mqk[Q16];
                     for (int j = 0; j < Q16; ++j) {