From: Johannes Gäßler <redacted>
Date: Sat, 23 Aug 2025 19:37:06 +0000 (+0200)
Subject: CUDA: fix half2 -> half conversion for HIP (llama/15529)
X-Git-Tag: v0.9.1~163
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=a61ecaac4791bf1672ddba7b3c6378b4a9187095;p=pkg%2Fggml%2Fsources%2Fggml

CUDA: fix half2 -> half conversion for HIP (llama/15529)
---

diff --git a/src/ggml-cuda/fattn-tile-f16.cu b/src/ggml-cuda/fattn-tile-f16.cu
index 6239d184..a900799a 100644
--- a/src/ggml-cuda/fattn-tile-f16.cu
+++ b/src/ggml-cuda/fattn-tile-f16.cu
@@ -258,7 +258,7 @@ static __global__ void flash_attn_tile_ext_f16(
             const half val = hexp(sink - kqmax[j0/nwarps]);
             kqsum[j0/nwarps] = kqsum[j0/nwarps] * KQ_max_scale;
             if (threadIdx.x == 0) {
-                kqsum[j0/nwarps].x = __hadd(kqsum[j0/nwarps].x, val);
+                kqsum[j0/nwarps].x = __hadd(__low2half(kqsum[j0/nwarps]), val);
             }
 
 #pragma unroll