Repository: ggml-org/llama.cpp
Pull request: HIP: enable mfma mmq on gfx908 and gfx90a for select datatypes and shapes

This PR enables the MFMA path graciously provided by @deepsek on CDNA1/gfx908 and CDNA2/gfx90a devices where this is more performant than the current blas path.

This PR is fairly careful and only enables the path on datatype and batch sizes where tested to be more performant. It is likely that other data types would also benefit, and i will follow up with more at a later time.

<details>

<summary>Measurements</summary>

All mesurents on gfx908, pp1024

## Performance change:

| Model            |   Microbatch size | Test   |   t/s master |   t/s mi100mfma |   Speedup |
|:-----------------|------------------:|:-------|-------------:|----------------:|----------:|
| llama 13B Q5_K_M |                32 | pp1024 |       141.71 |          385.04 |      2.72 |
| llama 13B Q5_K_M |                64 | pp1024 |       214.31 |          575.49 |      2.69 |
| llama 13B Q5_K_M |               128 | pp1024 |       388.50 |          707.50 |      1.82 |
| llama 13B Q5_K_M |               256 | pp1024 |       653.29 |          777.11 |      1.19 |
| llama 13B Q5_K_M |               512 | pp1024 |       850.12 |          847.13 |      1.00 |
| llama 13B Q5_K_M |              1024 | pp1024 |      1137.55 |         1136.94 |      1.00 |
| llama 13B Q8_0   |                32 | pp1024 |       159.84 |          441.99 |      2.77 |
| llama 13B Q8_0   |                64 | pp1024 |       263.17 |          627.86 |      2.39 |
| llama 13B Q8_0   |               128 | pp1024 |       485.83 |          637.20 |      1.31 |
| llama 13B Q8_0   |               256 | pp1024 |       749.96 |          763.62 |      1.02 |
| llama 13B Q8_0   |               512 | pp1024 |       915.36 |          934.04 |      1.02 |
| llama 13B Q8_0   |              1024 | pp1024 |      1232.74 |         1216.02 |      0.99 |
| llama 7B Q4_0    |                32 | pp1024 |       465.36 |         1227.27 |      2.64 |
| llama 7B Q4_0    |                64 | pp1024 |       728.69 |         1615.49 |      2.22 |
| llama 7B Q4_0    |               128 | pp1024 |      1105.51 |         1850.48 |      1.67 |
| llama 7B Q4_0    |               256 | pp1024 |      1553.94 |         2158.73 |      1.39 |
| llama 7B Q4_0    |               512 | pp1024 |      1890.85 |         2220.39 |      1.17 |
| llama 7B Q4_0    |              1024 | pp1024 |      2115.34 |         2291.17 |      1.08 |
| llama 7B Q5_0    |                32 | pp1024 |       416.92 |         1016.73 |      2.44 |
| llama 7B Q5_0    |                64 | pp1024 |       757.01 |         1449.28 |      1.91 |
| llama 7B Q5_0    |               128 | pp1024 |      1125.59 |         1738.29 |      1.54 |
| llama 7B Q5_0    |               256 | pp1024 |      1567.06 |         2025.97 |      1.29 |
| llama 7B Q5_0    |               512 | pp1024 |      1875.81 |         2094.18 |      1.12 |
| llama 7B Q5_0    |              1024 | pp1024 |      1978.71 |         2153.81 |      1.09 |
| llama 7B Q6_K    |                32 | pp1024 |       312.87 |          775.24 |      2.48 |
| llama 7B Q6_K    |                64 | pp1024 |       688.96 |         1026.81 |      1.49 |
| llama 7B Q6_K    |               128 | pp1024 |      1064.16 |         1248.84 |      1.17 |
| llama 7B Q6_K    |               256 | pp1024 |      1484.92 |         1470.56 |      0.99 |
| llama 7B Q6_K    |               512 | pp1024 |      1731.99 |         1719.56 |      0.99 |
| llama 7B Q6_K    |              1024 | pp1024 |      1836.89 |         1835.06 |      1.00 |
| llama 8B Q4_K_M  |                32 | pp1024 |       301.40 |         1023.17 |      3.39 |
| llama 8B Q4_K_M  |                64 | pp1024 |       628.02 |         1426.31 |      2.27 |
| llama 8B Q4_K_M  |               128 | pp1024 |      1197.32 |         1814.15 |      1.52 |
| llama 8B Q4_K_M  |               256 | pp1024 |      1986.34 |         2217.39 |      1.12 |
| llama 8B Q4_K_M  |               512 | pp1024 |      2498.98 |         2490.22 |      1.00 |
| llama 8B Q4_K_M  |              1024 | pp1024 |      2726.83 |         2718.23 |      1.00 |

## MFMA forced for all datatypes and batch sizes:

```
| model                          |     params | n_ubatch |                  t/s |
| ------------------------------ | ---------: | -------: | -------------------: |
| llama 7B Q4_0                  |     6.74 B |       32 |       1278.43 ± 3.66 |
| llama 7B Q4_0                  |     6.74 B |       64 |       1722.63 ± 4.00 |
| llama 7B Q4_0                  |     6.74 B |      128 |       2009.38 ± 5.06 |
| llama 7B Q4_0                  |     6.74 B |      256 |       2401.78 ± 2.37 |
| llama 7B Q4_0                  |     6.74 B |      512 |       2470.53 ± 7.83 |
| llama 7B Q4_0                  |     6.74 B |     1024 |       2571.22 ± 2.30 |
| llama 8B Q4_K - Medium         |     8.03 B |       32 |       1080.02 ± 2.33 |
| llama 8B Q4_K - Medium         |     8.03 B |       64 |       1553.58 ± 1.27 |
| llama 8B Q4_K - Medium         |     8.03 B |      128 |       2012.27 ± 0.61 |
| llama 8B Q4_K - Medium         |     8.03 B |      256 |       2286.34 ± 0.87 |
| llama 8B Q4_K - Medium         |     8.03 B |      512 |       2358.34 ± 1.46 |
| llama 8B Q4_K - Medium         |     8.03 B |     1024 |       2376.71 ± 0.72 |
| llama 7B Q5_0                  |     6.74 B |       32 |       1080.44 ± 1.48 |
| llama 7B Q5_0                  |     6.74 B |       64 |       1589.45 ± 0.79 |
| llama 7B Q5_0                  |     6.74 B |      128 |       1897.23 ± 4.16 |
| llama 7B Q5_0                  |     6.74 B |      256 |       2250.23 ± 1.01 |
| llama 7B Q5_0                  |     6.74 B |      512 |       2320.69 ± 1.50 |
| llama 7B Q5_0                  |     6.74 B |     1024 |       2342.06 ± 6.30 |
| llama 7B Q6_K                  |     6.74 B |       32 |        843.44 ± 2.82 |
| llama 7B Q6_K                  |     6.74 B |       64 |       1137.25 ± 2.91 |
| llama 7B Q6_K                  |     6.74 B |      128 |       1392.07 ± 0.57 |
| llama 7B Q6_K                  |     6.74 B |      256 |       1568.60 ± 0.38 |
| llama 7B Q6_K                  |     6.74 B |      512 |       1597.53 ± 0.40 |
| llama 7B Q6_K                  |     6.74 B |     1024 |       1620.83 ± 3.65 |
| llama 13B Q8_0                 |    23.57 B |       32 |        453.19 ± 0.44 |
| llama 13B Q8_0                 |    23.57 B |       64 |        674.65 ± 0.18 |
| llama 13B Q8_0                 |    23.57 B |      128 |        663.48 ± 3.70 |
| llama 13B Q8_0                 |    23.57 B |      256 |        739.34 ± 0.95 |
| llama 13B Q8_0                 |    23.57 B |      512 |        770.74 ± 1.21 |
| llama 13B Q8_0                 |    23.57 B |     1024 |        779.23 ± 1.25 |
```

</details>


Changed files:
filename: ggml/src/ggml-cuda/common.cuh
patch: @@ -227,9 +227,9 @@ typedef float2 dfloat2;
 #define FP16_MMA_AVAILABLE
 #endif // defined(GGML_HIP_ROCWMMA_FATTN) && (defined(CDNA) || defined(RDNA3) || (defined(GGML_HIP_ROCWMMA_FATTN_GFX12) && defined(RDNA4)))
 
-#if defined(GGML_USE_HIP) && defined(CDNA3) && !defined(GGML_HIP_NO_MMQ_MFMA)
+#if defined(GGML_USE_HIP) && defined(CDNA) && !defined(GGML_HIP_NO_MMQ_MFMA)
 #define AMD_MFMA_AVAILABLE
-#endif // defined(GGML_USE_HIP) && defined(CDNA3) && !defined(GGML_HIP_NO_MMQ_MFMA)
+#endif // defined(GGML_USE_HIP) && defined(CDNA) && !defined(GGML_HIP_NO_MMQ_MFMA)
 
 #if !defined(GGML_USE_HIP) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING
 #define NEW_MMA_AVAILABLE
@@ -293,10 +293,9 @@ static bool fp32_mma_hardware_available(const int cc) {
     return GGML_CUDA_CC_IS_CDNA(cc);
 }
 
-// AMD CDNA3 matrix cores.. Will add support for other CDNA generations later.
 static bool amd_mfma_available(const int cc) {
 #if !defined(GGML_HIP_NO_MMQ_MFMA)
-    return GGML_CUDA_CC_IS_CDNA3(cc);
+    return GGML_CUDA_CC_IS_CDNA(cc);
 #else
     return false;
 #endif //!defined(GGML_HIP_NO_MMQ_MFMA)

filename: ggml/src/ggml-cuda/mmq.cu
patch: @@ -109,8 +109,8 @@ void ggml_cuda_mul_mat_q(
     const int64_t s03 = src0->nb[3] / ts_src0;
     const int64_t s3  =  dst->nb[3] / ts_dst;
 
-    const bool use_stream_k = ((GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA)
-                            || (GGML_CUDA_CC_IS_AMD(cc) && GGML_CUDA_CC_IS_CDNA3(cc)));
+    const bool use_stream_k = (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA)
+                            || GGML_CUDA_CC_IS_CDNA(cc);
 
     if (!ids) {
         const size_t nbytes_src1_q8_1 = ne13*ne12 * ne11*ne10_padded * sizeof(block_q8_1)/QK8_1 +
@@ -252,7 +252,7 @@ void ggml_cuda_op_mul_mat_q(
     // Also its fixup needs to allocate a temporary buffer in the memory pool.
     // There are multiple parallel CUDA streams for src1_ncols != ne11 which would introduce a race condition for this buffer.
     const bool use_stream_k = ((GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA)
-                            || (GGML_CUDA_CC_IS_AMD(cc) && GGML_CUDA_CC_IS_CDNA3(cc)))
+                            || GGML_CUDA_CC_IS_CDNA(cc))
                             && src1_ncols == ne11;
     const mmq_args args = {
         src0_dd_i, src0->type, (const int *) src1_ddq_i, nullptr, nullptr, dst_dd_i,
@@ -306,7 +306,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) {
         return false;
     }
 
-    if (new_mma_available(cc) || amd_mfma_available(cc)) {
+    if (new_mma_available(cc)) {
         return true;
     }
 
@@ -322,5 +322,21 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) {
         return !fp16_mma_hardware_available(cc) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
     }
 
+    if (amd_mfma_available(cc)) {
+        // As of ROCM 7.0 rocblas/tensile performs very poorly on CDNA3 and hipblaslt (via ROCBLAS_USE_HIPBLASLT)
+        // performs better but is currently suffering from a crash on this architecture.
+        // TODO: Revisit when hipblaslt is fixed on CDNA3
+        if (GGML_CUDA_CC_IS_CDNA3(cc)) {
+            return true;
+        }
+        if (ne11 <= 128 || type == GGML_TYPE_Q4_0 || type == GGML_TYPE_Q4_1 || type == GGML_TYPE_Q5_0 || type == GGML_TYPE_Q5_1) {
+            return true;
+        }
+        if (ne11 <= 256 && (type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q5_K)) {
+            return true;
+        }
+        return false;
+    }
+
     return (!GGML_CUDA_CC_IS_RDNA4(cc) && !GGML_CUDA_CC_IS_RDNA3(cc) && !GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
 }

filename: ggml/src/ggml-cuda/mmq.cuh
patch: @@ -3096,8 +3096,8 @@ static __global__ void mul_mat_q(
     }
     __syncthreads();
 
-    // On AMD or old CUDA the performance with stream-k was worse, use conventional tiling instead:
-#if (defined(GGML_USE_HIP) && !defined(CDNA3)) || __CUDA_ARCH__ < GGML_CUDA_CC_VOLTA
+    // On non-CDNA AMD or old CUDA the performance with stream-k was worse, use conventional tiling instead:
+#if (defined(GGML_USE_HIP) && !defined(CDNA)) || __CUDA_ARCH__ < GGML_CUDA_CC_VOLTA
     {
         const int wt = blockIdx.z / nchannels_y;
         const int zt = blockIdx.z - wt*nchannels_y;