Repository: ggml-org/llama.cpp
Pull request: HIP: use v_dot2_f32_f16 instruction for FA

See https://github.com/iacopPBK/llama.cpp-gfx906 . The fork uses an instruction for FP16 multiply-add with FP32 accumulation. This PR adopts the same instruction for the tile FA kernel.

| GPU         | Model         | FlashAttention   |   Microbatch size | Test    |   t/s fe1c92cd7 |   t/s d91e76574 |   Speedup |
|:------------|:--------------|:-----------------|------------------:|:--------|----------------:|----------------:|----------:|
| MI60 / MI50 | gemma 2B Q4_0 | Yes              |                16 | pp16384 |          329.07 |          629.31 |      1.91 |
| MI60 / MI50 | gemma 2B Q4_0 | Yes              |                32 | pp16384 |          309.59 |          728.92 |      2.35 |
| MI60 / MI50 | gemma 2B Q4_0 | Yes              |               512 | pp16384 |          397.50 |         1412.22 |      3.55 |
| MI60 / MI50 | llama 1B Q4_0 | Yes              |                16 | pp16384 |          682.84 |          922.76 |      1.35 |
| MI60 / MI50 | llama 1B Q4_0 | Yes              |                32 | pp16384 |          953.88 |         1187.68 |      1.25 |
| MI60 / MI50 | llama 1B Q4_0 | Yes              |               512 | pp16384 |         1510.71 |         2278.62 |      1.51 |
| MI60 / MI50 | llama 8B Q4_0 | Yes              |                16 | pp16384 |          193.82 |          278.56 |      1.44 |
| MI60 / MI50 | llama 8B Q4_0 | Yes              |                32 | pp16384 |          163.34 |          334.36 |      2.05 |
| MI60 / MI50 | llama 8B Q4_0 | Yes              |               512 | pp16384 |          204.03 |          504.45 |      2.47 |


Changed files:
filename: ggml/src/ggml-cuda/common.cuh
patch: @@ -545,6 +545,31 @@ static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, i
 #endif // defined(GGML_USE_HIP)
 }
 
+static __device__ __forceinline__ void ggml_cuda_mad(float & acc, const float v, const float u) {
+    acc += v*u;
+}
+
+static __device__ __forceinline__ void ggml_cuda_mad(float & acc, const float2 v, const float2 u) {
+    acc += v.x*u.x;
+    acc += v.y*u.y;
+}
+
+static __device__ __forceinline__ void ggml_cuda_mad(float & acc, const half2 v, const half2 u) {
+#if defined(GGML_USE_HIP) && defined(GCN)
+    asm volatile("v_dot2_f32_f16 %0, %1, %2, %0" : "+v"(acc) : "v"(v), "v"(u));
+#else
+#ifdef FAST_FP16_AVAILABLE
+    const float2 tmp = __half22float2(v*u);
+    acc += tmp.x + tmp.y;
+#else
+    const float2 tmpv = __half22float2(v);
+    const float2 tmpu = __half22float2(u);
+    acc += tmpv.x * tmpu.x;
+    acc += tmpv.y * tmpu.y;
+#endif // FAST_FP16_AVAILABLE
+#endif // defined(GGML_USE_HIP) && defined(GCN)
+}
+
 static __device__ __forceinline__ float ggml_cuda_e8m0_to_fp32(uint8_t x) {
 #if CUDART_VERSION >= 12080
     const nv_bfloat16 e = __nv_cvt_e8m0_to_bf16raw(x);

filename: ggml/src/ggml-cuda/fattn-tile.cu
patch: @@ -304,12 +304,7 @@ static __global__ void flash_attn_tile(
                 for (int i_KQ_0 = 0; i_KQ_0 < kq_stride; i_KQ_0 += warp_size) {
 #pragma unroll
                     for (int j_KQ_0 = 0; j_KQ_0 < ncols; j_KQ_0 += nwarps) {
-#ifdef FAST_FP16_AVAILABLE
-                        const float2 tmp = __half22float2(K_k[i_KQ_0/warp_size] * Q_k[j_KQ_0/nwarps]);
-                        sum[i_KQ_0/warp_size][j_KQ_0/nwarps] += tmp.x + tmp.y;
-#else
-                        sum[i_KQ_0/warp_size][j_KQ_0/nwarps] += K_k[i_KQ_0/warp_size] * Q_k[j_KQ_0/nwarps];
-#endif // FAST_FP16_AVAILABLE
+                        ggml_cuda_mad(sum[i_KQ_0/warp_size][j_KQ_0/nwarps], K_k[i_KQ_0/warp_size], Q_k[j_KQ_0/nwarps]);
                     }
                 }
             }