CUDA: fix half2 -> half conversion for HIP (ggml-org#15529)

JohannesGaessler · web-flow · commit 710dfc465a68 · 2025-08-23T21:37:06.000+02:00
diff --git a/ggml/src/ggml-cuda/fattn-tile-f16.cu b/ggml/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

Original file line number	Diff line number	Diff line change
`@@ -258,7 +258,7 @@ static __global__ void flash_attn_tile_ext_f16(`
`258`	`258`	`const half val = hexp(sink - kqmax[j0/nwarps]);`
`259`	`259`	`kqsum[j0/nwarps] = kqsum[j0/nwarps] * KQ_max_scale;`
`260`	`260`	`if (threadIdx.x == 0) {`
`261`		`- kqsum[j0/nwarps].x = __hadd(kqsum[j0/nwarps].x, val);`
	`261`	`+ kqsum[j0/nwarps].x = __hadd(__low2half(kqsum[j0/nwarps]), val);`
`262`	`262`	`}`
`263`	`263`
`264`	`264`	`#pragma unroll`