Repository: ggml-org/llama.cpp
Pull request: HIP: enable WMMA-MMQ INT kernels for RDNA 3

Enabled WMMA-MMQ INT kernels for RDNA 3 architecture on AMD GPUs

Following similar approach to https://github.com/ggml-org/llama.cpp/pull/17156

Using ./build/bin/llama-bench to collect the following performance results

Build command for the following performance results:
HIPCXX="$(hipconfig -l)/clang" HIP_PATH="$(hipconfig -R)" cmake -S . -B build -DGGML_HIP=ON -DGGML_CUDA_FORCE_MMQ=OFF -DGGML_HIP_UMA=OFF -DGGML_HIP_ROCWMMA_FATTN=OFF -DGPU_TARGETS="gfx1100" -DGGML_HIP_GRAPHS=OFF -DLLAMA_CURL=OFF -DGGML_CUDA_FORCE_CUBLAS=OFF -DCMAKE_BUILD_TYPE=Release && cmake --build build --config Release -- -j 32

<details>
  <summary><b>Popular models performance result for AMD Radeon AI PRO W7900 (gfx1100)</b></summary>
<img width="1663" height="976" alt="image" src="https://github.com/user-attachments/assets/b53bc2f1-26a9-4ee0-beee-2027b10fd1d8" />
<img width="1668" height="660" alt="image" src="https://github.com/user-attachments/assets/9bc9737c-c6c5-411b-9e74-d49aad3f6209" />
</details>

<details>
  <summary><b>Popular models performance result for AMD Strix Halo (gfx1151)</b></summary>
<img width="2053" height="962" alt="image" src="https://github.com/user-attachments/assets/9256c14b-bd57-40be-9dc1-e666f0bf560e" />
<img width="1936" height="1070" alt="image" src="https://github.com/user-attachments/assets/daa8d3f2-ae35-4e85-a032-cd6cdefd07d4" />
</details>

<details>
  <summary><b>All quantization performance result for AMD Radeon AI PRO W7900 (gfx1100)</b></summary>

| GPU                 | Model                         |   Microbatch size | Test   |   t/s 583cb834 |   t/s mmq_feature_branch |   Speedup |
|:--------------------|:------------------------------|------------------:|:-------|---------------:|------------------------------------------------------:|----------:|
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |                 1 | pp2048 |         119.44 |                                                119.96 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |                 2 | pp2048 |         196.73 |                                                197.43 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |                 4 | pp2048 |         303.89 |                                                304.92 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |                 8 | pp2048 |         445.70 |                                                507.12 |      1.14 |
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |                16 | pp2048 |         634.63 |                                                804.82 |      1.27 |
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |                32 | pp2048 |         886.68 |                                               1027.74 |      1.16 |
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |                64 | pp2048 |         981.10 |                                               1445.47 |      1.47 |
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |               128 | pp2048 |        1510.05 |                                               1679.38 |      1.11 |
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |               256 | pp2048 |        2047.07 |                                               2016.22 |      0.98 |
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |               512 | pp2048 |        2012.02 |                                               2079.92 |      1.03 |
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |              1024 | pp2048 |        1945.28 |                                               2112.55 |      1.09 |
| PRO W7900 Dual Slot | llama 8B IQ1_S - 1.5625 bpw   |              2048 | pp2048 |        1976.88 |                                               2072.28 |      1.05 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |                 1 | pp2048 |          92.63 |                                                 92.72 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |                 2 | pp2048 |         156.95 |                                                156.95 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |                 4 | pp2048 |         254.64 |                                                254.45 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |                 8 | pp2048 |         383.33 |                                                426.76 |      1.11 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |                16 | pp2048 |         535.59 |                                                458.25 |      0.86 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |                32 | pp2048 |         783.94 |                                                906.32 |      1.16 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |                64 | pp2048 |         945.33 |                                               1373.64 |      1.45 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |               128 | pp2048 |        1490.37 |                                               1465.06 |      0.98 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |               256 | pp2048 |        2016.36 |                                               1728.85 |      0.86 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |               512 | pp2048 |        1986.67 |                                               1824.80 |      0.92 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |              1024 | pp2048 |        1929.90 |                                               1853.21 |      0.96 |
| PRO W7900 Dual Slot | llama 8B IQ2_S - 2.5 bpw      |              2048 | pp2048 |        1953.60 |                                               1808.53 |      0.93 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |                 1 | pp2048 |          95.74 |                                                 95.98 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |                 2 | pp2048 |         160.46 |                                                160.64 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |                 4 | pp2048 |         255.90 |                                                256.54 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |                 8 | pp2048 |         378.20 |                                                421.88 |      1.12 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |                16 | pp2048 |         547.54 |                                                453.34 |      0.83 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |                32 | pp2048 |         777.44 |                                                963.99 |      1.24 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |                64 | pp2048 |         948.91 |                                               1345.01 |      1.42 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |               128 | pp2048 |        1478.42 |                                               1429.39 |      0.97 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |               256 | pp2048 |        2027.36 |                                               1694.89 |      0.84 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |               512 | pp2048 |        1993.90 |                                               1788.80 |      0.90 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |              1024 | pp2048 |        1942.02 |                                               1827.91 |      0.94 |
| PRO W7900 Dual Slot | llama 8B IQ2_XS - 2.3125 bpw  |              2048 | pp2048 |        1974.40 |                                               1799.10 |      0.91 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |                 1 | pp2048 |          79.60 |                                                 79.68 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |                 2 | pp2048 |         139.88 |                                                139.72 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |                 4 | pp2048 |         235.12 |                                                235.01 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |                 8 | pp2048 |         350.25 |                                                385.55 |      1.10 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |                16 | pp2048 |         525.61 |                                                639.74 |      1.22 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |                32 | pp2048 |         761.19 |                                                671.28 |      0.88 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |                64 | pp2048 |         962.85 |                                               1338.24 |      1.39 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |               128 | pp2048 |        1492.21 |                                               1716.92 |      1.15 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |               256 | pp2048 |        2033.28 |                                               1989.56 |      0.98 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |               512 | pp2048 |        1997.97 |                                               2053.91 |      1.03 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |              1024 | pp2048 |        1940.79 |                                               2083.00 |      1.07 |
| PRO W7900 Dual Slot | llama 8B IQ2_XXS - 2.0625 bpw |              2048 | pp2048 |        1975.83 |                                               2040.61 |      1.03 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |                 1 | pp2048 |          75.52 |                                                 75.55 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |                 2 | pp2048 |         134.05 |                                                134.03 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |                 4 | pp2048 |         229.54 |                                                228.94 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |                 8 | pp2048 |         366.06 |                                                404.29 |      1.10 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |                16 | pp2048 |         494.08 |                                                526.54 |      1.07 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |                32 | pp2048 |         757.72 |                                                587.11 |      0.77 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |                64 | pp2048 |         908.26 |                                               1357.06 |      1.49 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |               128 | pp2048 |        1452.44 |                                               1707.32 |      1.18 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |               256 | pp2048 |        2006.58 |                                               2000.21 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |               512 | pp2048 |        1974.82 |                                               2058.28 |      1.04 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |              1024 | pp2048 |        1924.67 |                                               2075.83 |      1.08 |
| PRO W7900 Dual Slot | llama 8B IQ3_S - 3.4375 bpw   |              2048 | pp2048 |        1952.66 |                                               2011.09 |      1.03 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |                 1 | pp2048 |          74.91 |                                                 74.97 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |                 2 | pp2048 |         131.88 |                                                131.83 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |                 4 | pp2048 |         221.86 |                                                221.47 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |                 8 | pp2048 |         344.87 |                                                379.84 |      1.10 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |                16 | pp2048 |         504.66 |                                                554.38 |      1.10 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |                32 | pp2048 |         762.33 |                                                615.27 |      0.81 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |                64 | pp2048 |         915.10 |                                               1367.04 |      1.49 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |               128 | pp2048 |        1439.75 |                                               1701.67 |      1.18 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |               256 | pp2048 |        1990.07 |                                               2002.38 |      1.01 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |               512 | pp2048 |        1963.91 |                                               2056.56 |      1.05 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |              1024 | pp2048 |        1917.01 |                                               2072.95 |      1.08 |
| PRO W7900 Dual Slot | llama 8B IQ3_S mix - 3.66 bpw |              2048 | pp2048 |        1947.19 |                                               2016.53 |      1.04 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |                 1 | pp2048 |          82.83 |                                                 82.38 |      0.99 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |                 2 | pp2048 |         145.51 |                                                145.33 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |                 4 | pp2048 |         244.41 |                                                243.89 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |                 8 | pp2048 |         376.21 |                                                417.63 |      1.11 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |                16 | pp2048 |         525.57 |                                                593.25 |      1.13 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |                32 | pp2048 |         787.13 |                                                620.31 |      0.79 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |                64 | pp2048 |         909.63 |                                               1401.15 |      1.54 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |               128 | pp2048 |        1441.25 |                                               1717.14 |      1.19 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |               256 | pp2048 |        1992.22 |                                               1967.66 |      0.99 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |               512 | pp2048 |        1966.36 |                                               2002.64 |      1.02 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |              1024 | pp2048 |        1921.96 |                                               1950.32 |      1.01 |
| PRO W7900 Dual Slot | llama 8B IQ3_XS - 3.3 bpw     |              2048 | pp2048 |        1952.51 |                                               1808.72 |      0.93 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |                 1 | pp2048 |          89.59 |                                                 89.49 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |                 2 | pp2048 |         154.29 |                                                154.07 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |                 4 | pp2048 |         253.94 |                                                253.34 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |                 8 | pp2048 |         376.96 |                                                418.83 |      1.11 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |                16 | pp2048 |         534.30 |                                                608.85 |      1.14 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |                32 | pp2048 |         795.20 |                                                672.18 |      0.85 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |                64 | pp2048 |         907.68 |                                               1411.46 |      1.56 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |               128 | pp2048 |        1435.77 |                                               1727.08 |      1.20 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |               256 | pp2048 |        1980.00 |                                               2039.01 |      1.03 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |               512 | pp2048 |        1959.33 |                                               2103.84 |      1.07 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |              1024 | pp2048 |        1916.08 |                                               2120.41 |      1.11 |
| PRO W7900 Dual Slot | llama 8B IQ3_XXS - 3.0625 bpw |              2048 | pp2048 |        1952.05 |                                               2058.58 |      1.05 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |                 1 | pp2048 |          89.84 |                                                 90.14 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |                 2 | pp2048 |         161.96 |                                                162.00 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |                 4 | pp2048 |         283.82 |                                                284.01 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |                 8 | pp2048 |         457.70 |                                                518.59 |      1.13 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |                16 | pp2048 |         621.81 |                                                795.53 |      1.28 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |                32 | pp2048 |         878.06 |                                                821.83 |      0.94 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |                64 | pp2048 |         878.89 |                                               1558.47 |      1.77 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |               128 | pp2048 |        1414.02 |                                               1858.45 |      1.31 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |               256 | pp2048 |        1975.53 |                                               2192.85 |      1.11 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |               512 | pp2048 |        1951.61 |                                               2243.64 |      1.15 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |              1024 | pp2048 |        1923.83 |                                               2260.91 |      1.18 |
| PRO W7900 Dual Slot | llama 8B IQ4_NL - 4.5 bpw     |              2048 | pp2048 |        1977.10 |                                               2204.06 |      1.11 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |                 1 | pp2048 |          94.01 |                                                 94.10 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |                 2 | pp2048 |         170.10 |                                                170.09 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |                 4 | pp2048 |         298.65 |                                                298.65 |      1.00 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |                 8 | pp2048 |         466.09 |                                                532.42 |      1.14 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |                16 | pp2048 |         634.03 |                                                824.04 |      1.30 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |                32 | pp2048 |         886.40 |                                                597.30 |      0.67 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |                64 | pp2048 |         888.15 |                                               1574.37 |      1.77 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |               128 | pp2048 |        1432.87 |                                               1865.24 |      1.30 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |               256 | pp2048 |        1983.10 |                                               2191.29 |      1.10 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |               512 | pp2048 |        1957.77 |                                               2237.67 |      1.14 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |              1024 | pp2048 |        1927.76 |                                               2258.80 |      1.17 |
| PRO W7900 Dual Slot | llama 8B IQ4_XS - 4.25 bpw    |              2048 | pp2048 |        1966.60 |                                               2200.45 |      1.12 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |                 1 | pp2048 |         104.03 |                                                103.86 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |                 2 | pp2048 |         156.56 |                                                156.27 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |                 4 | pp2048 |         204.91 |                                                204.39 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |                 8 | pp2048 |         249.14 |                                                265.99 |      1.07 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |                16 | pp2048 |         557.05 |                                                409.44 |      0.74 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |                32 | pp2048 |         549.89 |                                                579.28 |      1.05 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |                64 | pp2048 |         921.02 |                                                818.31 |      0.89 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |               128 | pp2048 |        1452.22 |                                               1164.26 |      0.80 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |               256 | pp2048 |        1991.42 |                                               1221.65 |      0.61 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |               512 | pp2048 |        1985.17 |                                               1333.52 |      0.67 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |              1024 | pp2048 |        1942.76 |                                               1435.10 |      0.74 |
| PRO W7900 Dual Slot | llama 8B Q2_K_S               |              2048 | pp2048 |        1975.16 |                                               1434.72 |      0.73 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |                 1 | pp2048 |          80.44 |                                                 80.66 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |                 2 | pp2048 |         134.57 |                                                134.67 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |                 4 | pp2048 |         197.14 |                                                197.57 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |                 8 | pp2048 |         246.28 |                                                264.20 |      1.07 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |                16 | pp2048 |         539.32 |                                                469.68 |      0.87 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |                32 | pp2048 |         756.66 |                                               1011.50 |      1.34 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |                64 | pp2048 |         832.07 |                                               1378.54 |      1.66 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |               128 | pp2048 |        1331.93 |                                               1660.21 |      1.25 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |               256 | pp2048 |        1834.19 |                                               1937.16 |      1.06 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |               512 | pp2048 |        1932.29 |                                               1994.61 |      1.03 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |              1024 | pp2048 |        1930.39 |                                               2013.32 |      1.04 |
| PRO W7900 Dual Slot | llama 8B Q3_K_S               |              2048 | pp2048 |        1971.61 |                                               1971.32 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |                 1 | pp2048 |          92.59 |                                                 92.85 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |                 2 | pp2048 |         167.66 |                                                168.00 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |                 4 | pp2048 |         293.21 |                                                293.11 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |                 8 | pp2048 |         477.35 |                                                545.22 |      1.14 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |                16 | pp2048 |         623.04 |                                                764.53 |      1.23 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |                32 | pp2048 |         906.53 |                                                640.59 |      0.71 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |                64 | pp2048 |         895.68 |                                               1526.66 |      1.70 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |               128 | pp2048 |        1437.32 |                                               1862.50 |      1.30 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |               256 | pp2048 |        2009.97 |                                               2178.61 |      1.08 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |               512 | pp2048 |        1979.84 |                                               2221.04 |      1.12 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |              1024 | pp2048 |        1933.14 |                                               2240.34 |      1.16 |
| PRO W7900 Dual Slot | llama 8B Q4_0                 |              2048 | pp2048 |        1990.05 |                                               2178.16 |      1.09 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |                 1 | pp2048 |          87.30 |                                                 87.02 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |                 2 | pp2048 |         159.37 |                                                158.97 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |                 4 | pp2048 |         274.23 |                                                274.11 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |                 8 | pp2048 |         478.54 |                                                545.67 |      1.14 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |                16 | pp2048 |         611.39 |                                                807.16 |      1.32 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |                32 | pp2048 |         927.59 |                                               1118.95 |      1.21 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |                64 | pp2048 |         877.06 |                                               1502.57 |      1.71 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |               128 | pp2048 |        1416.30 |                                               1653.30 |      1.17 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |               256 | pp2048 |        1989.90 |                                               2017.58 |      1.01 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |               512 | pp2048 |        1965.18 |                                               2079.99 |      1.06 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |              1024 | pp2048 |        1923.68 |                                               2116.16 |      1.10 |
| PRO W7900 Dual Slot | llama 8B Q4_1                 |              2048 | pp2048 |        1982.31 |                                               2079.10 |      1.05 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |                 1 | pp2048 |          78.77 |                                                 78.62 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |                 2 | pp2048 |         126.64 |                                                126.48 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |                 4 | pp2048 |         188.75 |                                                188.74 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |                 8 | pp2048 |         252.05 |                                                270.04 |      1.07 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |                16 | pp2048 |         590.43 |                                                812.30 |      1.38 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |                32 | pp2048 |         807.23 |                                               1085.64 |      1.34 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |                64 | pp2048 |         870.47 |                                               1480.21 |      1.70 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |               128 | pp2048 |        1396.57 |                                               1673.81 |      1.20 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |               256 | pp2048 |        1931.75 |                                               2052.97 |      1.06 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |               512 | pp2048 |        1957.64 |                                               2109.49 |      1.08 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |              1024 | pp2048 |        1931.27 |                                               2142.37 |      1.11 |
| PRO W7900 Dual Slot | llama 8B Q4_K_S               |              2048 | pp2048 |        1978.30 |                                               2106.29 |      1.06 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |                 1 | pp2048 |          70.60 |                                                 70.69 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |                 2 | pp2048 |         128.57 |                                                128.67 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |                 4 | pp2048 |         231.59 |                                                231.96 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |                 8 | pp2048 |         428.61 |                                                482.33 |      1.13 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |                16 | pp2048 |         495.18 |                                                578.44 |      1.17 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |                32 | pp2048 |         768.89 |                                                874.56 |      1.14 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |                64 | pp2048 |         757.85 |                                               1286.75 |      1.70 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |               128 | pp2048 |        1237.01 |                                               1487.90 |      1.20 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |               256 | pp2048 |        1733.45 |                                               1885.56 |      1.09 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |               512 | pp2048 |        1873.29 |                                               1974.93 |      1.05 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |              1024 | pp2048 |        1907.48 |                                               2027.97 |      1.06 |
| PRO W7900 Dual Slot | llama 8B Q5_1                 |              2048 | pp2048 |        1960.31 |                                               1995.51 |      1.02 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |                 1 | pp2048 |          73.36 |                                                 73.33 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |                 2 | pp2048 |         121.67 |                                                121.67 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |                 4 | pp2048 |         184.13 |                                                183.70 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |                 8 | pp2048 |         246.71 |                                                263.30 |      1.07 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |                16 | pp2048 |         554.75 |                                                829.08 |      1.49 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |                32 | pp2048 |         699.59 |                                               1116.24 |      1.60 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |                64 | pp2048 |         818.27 |                                               1498.33 |      1.83 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |               128 | pp2048 |        1332.58 |                                               1616.54 |      1.21 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |               256 | pp2048 |        1843.06 |                                               1968.65 |      1.07 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |               512 | pp2048 |        1928.32 |                                               2027.76 |      1.05 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |              1024 | pp2048 |        1934.12 |                                               2064.67 |      1.07 |
| PRO W7900 Dual Slot | llama 8B Q5_K_S               |              2048 | pp2048 |        1971.71 |                                               2026.68 |      1.03 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |                 1 | pp2048 |          69.29 |                                                 69.18 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |                 2 | pp2048 |         123.13 |                                                123.00 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |                 4 | pp2048 |         197.96 |                                                197.34 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |                 8 | pp2048 |         286.01 |                                                306.88 |      1.07 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |                16 | pp2048 |         498.59 |                                                623.24 |      1.25 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |                32 | pp2048 |         648.44 |                                                813.97 |      1.26 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |                64 | pp2048 |         834.21 |                                               1084.13 |      1.30 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |               128 | pp2048 |        1364.80 |                                               1094.80 |      0.80 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |               256 | pp2048 |        1897.44 |                                               1322.89 |      0.70 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |               512 | pp2048 |        1930.36 |                                               1450.17 |      0.75 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |              1024 | pp2048 |        1918.10 |                                               1470.46 |      0.77 |
| PRO W7900 Dual Slot | llama 8B Q6_K                 |              2048 | pp2048 |        1973.29 |                                               1488.04 |      0.75 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |                 1 | pp2048 |          62.95 |                                                 63.04 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |                 2 | pp2048 |         115.78 |                                                115.89 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |                 4 | pp2048 |         210.63 |                                                210.77 |      1.00 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |                 8 | pp2048 |         383.93 |                                                422.00 |      1.10 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |                16 | pp2048 |         576.38 |                                                717.52 |      1.24 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |                32 | pp2048 |         864.31 |                                                619.38 |      0.72 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |                64 | pp2048 |         789.64 |                                               1476.07 |      1.87 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |               128 | pp2048 |        1295.11 |                                               1815.55 |      1.40 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |               256 | pp2048 |        1874.73 |                                               2173.67 |      1.16 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |               512 | pp2048 |        1891.10 |                                               2227.88 |      1.18 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |              1024 | pp2048 |        1902.79 |                                               2251.06 |      1.18 |
| PRO W7900 Dual Slot | llama 8B Q8_0                 |              2048 | pp2048 |        1964.10 |                                               2197.58 |      1.12 |
</details>

<details>
  <summary><b>All quantization performance result for AMD Strix Halo (gfx1151)</b></summary>

| GPU      | Model                         |   Microbatch size | Test   |   t/s master |   t/s mmq_feature_branch |   Speedup |
|:---------|:------------------------------|------------------:|:-------|-------------:|------------------------------------------------------:|----------:|
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |                 1 | pp2048 |        59.83 |                                                 59.90 |      1.00 |
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |                 2 | pp2048 |       104.71 |                                                104.66 |      1.00 |
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |                 4 | pp2048 |       163.47 |                                                163.36 |      1.00 |
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |                 8 | pp2048 |       217.52 |                                                269.92 |      1.24 |
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |                16 | pp2048 |       389.90 |                                                537.71 |      1.38 |
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |                32 | pp2048 |       532.61 |                                                600.36 |      1.13 |
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |                64 | pp2048 |       252.24 |                                                770.75 |      3.06 |
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |               128 | pp2048 |       436.51 |                                                854.21 |      1.96 |
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |               256 | pp2048 |       604.08 |                                                914.89 |      1.51 |
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |               512 | pp2048 |       730.25 |                                                891.78 |      1.22 |
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |              1024 | pp2048 |       808.51 |                                                888.27 |      1.10 |
| Graphics | llama 8B IQ1_S - 1.5625 bpw   |              2048 | pp2048 |       784.05 |                                                807.68 |      1.03 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |                 1 | pp2048 |        49.70 |                                                 49.49 |      1.00 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |                 2 | pp2048 |        88.14 |                                                 87.73 |      1.00 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |                 4 | pp2048 |       139.02 |                                                138.65 |      1.00 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |                 8 | pp2048 |       178.29 |                                                212.21 |      1.19 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |                16 | pp2048 |       324.40 |                                                304.76 |      0.94 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |                32 | pp2048 |       456.98 |                                                550.06 |      1.20 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |                64 | pp2048 |       250.13 |                                                703.03 |      2.81 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |               128 | pp2048 |       433.34 |                                                737.41 |      1.70 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |               256 | pp2048 |       602.11 |                                                789.88 |      1.31 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |               512 | pp2048 |       727.09 |                                                781.97 |      1.08 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |              1024 | pp2048 |       806.37 |                                                782.58 |      0.97 |
| Graphics | llama 8B IQ2_XS - 2.3125 bpw  |              2048 | pp2048 |       782.47 |                                                724.99 |      0.93 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |                 1 | pp2048 |        41.54 |                                                 41.38 |      1.00 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |                 2 | pp2048 |        75.11 |                                                 74.86 |      1.00 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |                 4 | pp2048 |       123.87 |                                                123.62 |      1.00 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |                 8 | pp2048 |       164.17 |                                                191.86 |      1.17 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |                16 | pp2048 |       313.88 |                                                380.15 |      1.21 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |                32 | pp2048 |       457.55 |                                                419.77 |      0.92 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |                64 | pp2048 |       251.35 |                                                715.39 |      2.85 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |               128 | pp2048 |       435.15 |                                                832.72 |      1.91 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |               256 | pp2048 |       605.74 |                                                895.57 |      1.48 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |               512 | pp2048 |       714.66 |                                                876.64 |      1.23 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |              1024 | pp2048 |       807.75 |                                                871.60 |      1.08 |
| Graphics | llama 8B IQ2_XXS - 2.0625 bpw |              2048 | pp2048 |       781.91 |                                                798.25 |      1.02 |
| Graphics | llama 8B Q2_K_S               |                 1 | pp2048 |        51.50 |                                                 51.49 |      1.00 |
| Graphics | llama 8B Q2_K_S               |                 2 | pp2048 |        82.63 |                                                 82.75 |      1.00 |
| Graphics | llama 8B Q2_K_S               |                 4 | pp2048 |       106.98 |                                                107.28 |      1.00 |
| Graphics | llama 8B Q2_K_S               |                 8 | pp2048 |       115.64 |                                                129.42 |      1.12 |
| Graphics | llama 8B Q2_K_S               |                16 | pp2048 |       336.98 |                                                257.13 |      0.76 |
| Graphics | llama 8B Q2_K_S               |                32 | pp2048 |       323.61 |                                                360.17 |      1.11 |
| Graphics | llama 8B Q2_K_S               |                64 | pp2048 |       249.79 |                                                494.88 |      1.98 |
| Graphics | llama 8B Q2_K_S               |               128 | pp2048 |       432.75 |                                                544.20 |      1.26 |
| Graphics | llama 8B Q2_K_S               |               256 | pp2048 |       598.52 |                                                585.13 |      0.98 |
| Graphics | llama 8B Q2_K_S               |               512 | pp2048 |       729.62 |                                                602.31 |      0.83 |
| Graphics | llama 8B Q2_K_S               |              1024 | pp2048 |       800.13 |                                                627.29 |      0.78 |
| Graphics | llama 8B Q2_K_S               |              2048 | pp2048 |       778.95 |                                                591.66 |      0.76 |
| Graphics | llama 8B Q3_K_S               |                 1 | pp2048 |        42.14 |                                                 42.14 |      1.00 |
| Graphics | llama 8B Q3_K_S               |                 2 | pp2048 |        72.14 |                                                 71.93 |      1.00 |
| Graphics | llama 8B Q3_K_S               |                 4 | pp2048 |       103.72 |                                                103.01 |      0.99 |
| Graphics | llama 8B Q3_K_S               |                 8 | pp2048 |       115.52 |                                                128.04 |      1.11 |
| Graphics | llama 8B Q3_K_S               |                16 | pp2048 |       326.11 |                                                313.91 |      0.96 |
| Graphics | llama 8B Q3_K_S               |                32 | pp2048 |       442.67 |                                                589.17 |      1.33 |
| Graphics | llama 8B Q3_K_S               |                64 | pp2048 |       243.59 |                                                736.42 |      3.02 |
| Graphics | llama 8B Q3_K_S               |               128 | pp2048 |       423.89 |                                                864.64 |      2.04 |
| Graphics | llama 8B Q3_K_S               |               256 | pp2048 |       590.88 |                                                923.04 |      1.56 |
| Graphics | llama 8B Q3_K_S               |               512 | pp2048 |       717.73 |                                                912.07 |      1.27 |
| Graphics | llama 8B Q3_K_S               |              1024 | pp2048 |       799.04 |                                                897.25 |      1.12 |
| Graphics | llama 8B Q3_K_S               |              2048 | pp2048 |       775.37 |                                                812.76 |      1.05 |
| Graphics | llama 8B Q4_0                 |                 1 | pp2048 |        39.70 |                                                 39.79 |      1.00 |
| Graphics | llama 8B Q4_0                 |                 2 | pp2048 |        76.18 |                                                 76.37 |      1.00 |
| Graphics | llama 8B Q4_0                 |                 4 | pp2048 |       139.13 |                                                139.35 |      1.00 |
| Graphics | llama 8B Q4_0                 |                 8 | pp2048 |       226.14 |                                                283.05 |      1.25 |
| Graphics | llama 8B Q4_0                 |                16 | pp2048 |       374.19 |                                                494.95 |      1.32 |
| Graphics | llama 8B Q4_0                 |                32 | pp2048 |       509.56 |                                                379.15 |      0.74 |
| Graphics | llama 8B Q4_0                 |                64 | pp2048 |       243.16 |                                                808.46 |      3.32 |
| Graphics | llama 8B Q4_0                 |               128 | pp2048 |       423.11 |                                                916.38 |      2.17 |
| Graphics | llama 8B Q4_0                 |               256 | pp2048 |       590.42 |                                                979.33 |      1.66 |
| Graphics | llama 8B Q4_0                 |               512 | pp2048 |       726.45 |                                                955.42 |      1.32 |
| Graphics | llama 8B Q4_0                 |              1024 | pp2048 |       808.57 |                                                943.67 |      1.17 |
| Graphics | llama 8B Q4_0                 |              2048 | pp2048 |       779.26 |                                                855.14 |      1.10 |
| Graphics | llama 8B Q4_1                 |                 1 | pp2048 |        36.48 |                                                 36.52 |      1.00 |
| Graphics | llama 8B Q4_1                 |                 2 | pp2048 |        71.08 |                                                 71.24 |      1.00 |
| Graphics | llama 8B Q4_1                 |                 4 | pp2048 |       131.92 |                                                132.17 |      1.00 |
| Graphics | llama 8B Q4_1                 |                 8 | pp2048 |       219.34 |                                                274.38 |      1.25 |
| Graphics | llama 8B Q4_1                 |                16 | pp2048 |       341.72 |                                                486.38 |      1.42 |
| Graphics | llama 8B Q4_1                 |                32 | pp2048 |       491.51 |                                                623.27 |      1.27 |
| Graphics | llama 8B Q4_1                 |                64 | pp2048 |       241.80 |                                                791.80 |      3.27 |
| Graphics | llama 8B Q4_1                 |               128 | pp2048 |       420.49 |                                                815.73 |      1.94 |
| Graphics | llama 8B Q4_1                 |               256 | pp2048 |       589.14 |                                                880.82 |      1.50 |
| Graphics | llama 8B Q4_1                 |               512 | pp2048 |       722.50 |                                                872.68 |      1.21 |
| Graphics | llama 8B Q4_1                 |              1024 | pp2048 |       804.07 |                                                867.01 |      1.08 |
| Graphics | llama 8B Q4_1                 |              2048 | pp2048 |       776.24 |                                                799.71 |      1.03 |
| Graphics | llama 8B Q4_K_S               |                 1 | pp2048 |        36.85 |                                                 36.88 |      1.00 |
| Graphics | llama 8B Q4_K_S               |                 2 | pp2048 |        62.35 |                                                 62.45 |      1.00 |
| Graphics | llama 8B Q4_K_S               |                 4 | pp2048 |        94.11 |                                                 94.07 |      1.00 |
| Graphics | llama 8B Q4_K_S               |                 8 | pp2048 |       118.60 |                                                133.25 |      1.12 |
| Graphics | llama 8B Q5_1                 |                 1 | pp2048 |        30.45 |                                                 30.47 |      1.00 |
| Graphics | llama 8B Q5_1                 |                 2 | pp2048 |        59.52 |                                                 59.56 |      1.00 |
| Graphics | llama 8B Q5_1                 |                 4 | pp2048 |       112.25 |                                                112.25 |      1.00 |
| Graphics | llama 8B Q5_1                 |                 8 | pp2048 |       191.34 |                                                232.75 |      1.22 |
| Graphics | llama 8B Q5_1                 |                16 | pp2048 |       267.52 |                                                342.23 |      1.28 |
| Graphics | llama 8B Q5_1                 |                32 | pp2048 |       447.40 |                                                506.85 |      1.13 |
| Graphics | llama 8B Q5_1                 |                64 | pp2048 |       228.36 |                                                696.09 |      3.05 |
| Graphics | llama 8B Q5_1                 |               128 | pp2048 |       400.89 |                                                813.09 |      2.03 |
| Graphics | llama 8B Q5_1                 |               256 | pp2048 |       566.88 |                                                884.65 |      1.56 |
| Graphics | llama 8B Q5_1                 |               512 | pp2048 |       705.10 |                                                880.81 |      1.25 |
| Graphics | llama 8B Q5_1                 |              1024 | pp2048 |       793.35 |                                                876.75 |      1.11 |
| Graphics | llama 8B Q5_1                 |              2048 | pp2048 |       769.48 |                                                803.44 |      1.04 |
| Graphics | llama 8B Q5_K_S               |                 1 | pp2048 |        33.36 |                                                 33.34 |      1.00 |
| Graphics | llama 8B Q5_K_S               |                 2 | pp2048 |        58.02 |                                                 57.86 |      1.00 |
| Graphics | llama 8B Q5_K_S               |                 4 | pp2048 |        90.37 |                                                 90.15 |      1.00 |
| Graphics | llama 8B Q5_K_S               |                 8 | pp2048 |       115.17 |                                                128.41 |      1.11 |
| Graphics | llama 8B Q5_K_S               |                16 | pp2048 |       328.82 |                                                508.01 |      1.54 |
| Graphics | llama 8B Q5_K_S               |                32 | pp2048 |       431.21 |                                                615.53 |      1.43 |
| Graphics | llama 8B Q5_K_S               |                64 | pp2048 |       235.29 |                                                792.11 |      3.37 |
| Graphics | llama 8B Q5_K_S               |               128 | pp2048 |       411.76 |                                                864.48 |      2.10 |
| Graphics | llama 8B Q5_K_S               |               256 | pp2048 |       578.32 |                                                919.98 |      1.59 |
| Graphics | llama 8B Q5_K_S               |               512 | pp2048 |       717.19 |                                                908.41 |      1.27 |
| Graphics | llama 8B Q5_K_S               |              1024 | pp2048 |       799.98 |                                                900.27 |      1.13 |
| Graphics | llama 8B Q5_K_S               |              2048 | pp2048 |       778.09 |                                                820.54 |      1.05 |
| Graphics | llama 8B Q6_K                 |                 1 | pp2048 |        29.58 |                                                 29.51 |      1.00 |
| Graphics | llama 8B Q6_K                 |                 2 | pp2048 |        56.47 |                                                 56.37 |      1.00 |
| Graphics | llama 8B Q6_K                 |                 4 | pp2048 |       102.88 |                                                102.60 |      1.00 |
| Graphics | llama 8B Q6_K                 |                 8 | pp2048 |       143.89 |                                                163.08 |      1.13 |
| Graphics | llama 8B Q6_K                 |                16 | pp2048 |       294.77 |                                                383.68 |      1.30 |
| Graphics | llama 8B Q6_K                 |                32 | pp2048 |       392.14 |                                                469.49 |      1.20 |
| Graphics | llama 8B Q6_K                 |                64 | pp2048 |       238.05 |                                                592.30 |      2.49 |
| Graphics | llama 8B Q6_K                 |               128 | pp2048 |       415.60 |                                                620.97 |      1.49 |
| Graphics | llama 8B Q6_K                 |               256 | pp2048 |       582.53 |                                                658.80 |      1.13 |
| Graphics | llama 8B Q6_K                 |               512 | pp2048 |       719.82 |                                                654.72 |      0.91 |
| Graphics | llama 8B Q6_K                 |              1024 | pp2048 |       796.31 |                                                663.25 |      0.83 |
| Graphics | llama 8B Q6_K                 |              2048 | pp2048 |       773.72 |                                                624.57 |      0.81 |
| Graphics | llama 8B Q8_0                 |                 1 | pp2048 |        25.02 |                                                 25.00 |      1.00 |
| Graphics | llama 8B Q8_0                 |                 2 | pp2048 |        48.78 |                                                 48.77 |      1.00 |
| Graphics | llama 8B Q8_0                 |                 4 | pp2048 |        92.37 |                                                 92.37 |      1.00 |
| Graphics | llama 8B Q8_0                 |                 8 | pp2048 |       166.97 |                                                198.40 |      1.19 |
| Graphics | llama 8B Q8_0                 |                16 | pp2048 |       303.03 |                                                364.41 |      1.20 |
| Graphics | llama 8B Q8_0                 |                32 | pp2048 |       484.44 |                                                374.38 |      0.77 |
| Graphics | llama 8B Q8_0                 |                64 | pp2048 |       227.11 |                                                750.26 |      3.30 |
| Graphics | llama 8B Q8_0                 |               128 | pp2048 |       398.25 |                                                873.07 |      2.19 |
| Graphics | llama 8B Q8_0                 |               256 | pp2048 |       564.91 |                                                941.85 |      1.67 |
| Graphics | llama 8B Q8_0                 |               512 | pp2048 |       705.53 |                                                925.05 |      1.31 |
| Graphics | llama 8B Q8_0                 |              1024 | pp2048 |       786.56 |                                                915.31 |      1.16 |
| Graphics | llama 8B Q8_0                 |              2048 | pp2048 |       773.68 |                                                833.31 |      1.08 |
</details>


Changed files:
filename: ggml/src/ggml-cuda/common.cuh
patch: @@ -226,7 +226,7 @@ static const char * cu_get_error_str(CUresult err) {
 #define AMD_MFMA_AVAILABLE
 #endif // defined(GGML_USE_HIP) && defined(CDNA) && !defined(GGML_HIP_NO_MMQ_MFMA)
 
-#if defined(GGML_USE_HIP) && defined(RDNA4)
+#if defined(GGML_USE_HIP) && (defined(RDNA4) || defined(RDNA3))
 #define AMD_WMMA_AVAILABLE
 #endif // defined(GGML_USE_HIP) && defined(RDNA4)
 
@@ -294,7 +294,7 @@ static bool amd_mfma_available(const int cc) {
 }
 
 static bool amd_wmma_available(const int cc) {
-    return GGML_CUDA_CC_IS_RDNA4(cc);
+    return (GGML_CUDA_CC_IS_RDNA4(cc) || GGML_CUDA_CC_IS_RDNA3(cc));
 }
 
 static bool volta_mma_available(const int cc) {

filename: ggml/src/ggml-cuda/mma.cuh
patch: @@ -173,6 +173,9 @@ namespace ggml_cuda_mma {
 #elif defined(AMD_WMMA_AVAILABLE)
 #if defined(RDNA4)
         static constexpr int ne = I * J / 32;
+#elif defined(RDNA3)
+        static constexpr int ne = (I == 16 && J == 16) ? I * J / 32 : I * J / 16;
+#endif // defined(RDNA4)
         T x[ne] = {0};
 
         static constexpr __device__ bool supported() {
@@ -182,7 +185,11 @@ namespace ggml_cuda_mma {
 
         static __device__ __forceinline__ int get_i(const int l) {
             if constexpr (I == 16 && J == 16) {
+#if defined(RDNA4)
                 return 8 * (threadIdx.x / 16) + l;
+#elif defined(RDNA3)
+                return 2 * l + (threadIdx.x / 16);
+#endif // defined(RDNA4)
             } else {
                 NO_DEVICE_CODE;
                 return -1;
@@ -197,7 +204,6 @@ namespace ggml_cuda_mma {
                 return -1;
             }
         }
-#endif
 #else
         static constexpr int ne = I * J / 32;
         T x[ne] = {0};
@@ -284,6 +290,7 @@ namespace ggml_cuda_mma {
             }
         }
 #elif defined(AMD_WMMA_AVAILABLE)
+
         static constexpr int ne = I * J / 32;
         half2 x[ne] = {{0.0f, 0.0f}};
 
@@ -544,23 +551,40 @@ namespace ggml_cuda_mma {
         } else if constexpr (std::is_same_v<T, int>) {
             if constexpr (I == 16 && J == 4) {
                 int64_t * xi = (int64_t *) t.x;
+#if defined(RDNA4)
                 const int64_t * xs = (int64_t *) ((const int *) xs0 + (threadIdx.x % t.I) * stride + 2 * (threadIdx.x / t.I));
                 xi[0] = xs[0];
-
+#elif defined(RDNA3)
+                static_assert(tile<I,J,T>::ne >= 4, "fragment too small");
+                const int64_t * xs = (int64_t *) ((const int *) xs0 + (threadIdx.x % t.I) * stride);
+                xi[0] = xs[0];
+                xi[1] = xs[1];
+#endif // defined(RDNA4)
             }else if constexpr (I == 16 && J == 8) {
                 int64_t * xi = (int64_t *) t.x;
+#if defined(RDNA4)
                 const int64_t * xs = (int64_t *) ((const int *) xs0 + (threadIdx.x % t.I) * stride + 4 * (threadIdx.x / t.I));
                 xi[0] = xs[0];
 
                 const int64_t * xs1 = (int64_t *) ((const int *) xs0 + (threadIdx.x % t.I) * stride + 4 * (threadIdx.x / t.I) + 2);
                 xi[1] = xs1[0];
+#elif defined(RDNA3)
+                static_assert(tile<I,J,T>::ne >= 8, "fragment too small");
+                const int64_t * xs = (int64_t *) ((const int *) xs0 + (threadIdx.x % t.I) * stride);
+                // contiguous four 64-bit chunks per lane for the wider RDNA3 fragment
+                xi[0] = xs[0];
+                xi[1] = xs[1];
+                const int64_t * xs1 = xs + 2;
+                xi[2] = xs1[0];
+                xi[3] = xs1[1];
 
             }else{
                 NO_DEVICE_CODE;
             }
         } else {
             NO_DEVICE_CODE;
         }
+#endif // defined(RDNA4)
 #else
 #pragma unroll
         for (int l = 0; l < t.ne; ++l) {
@@ -858,12 +882,14 @@ namespace ggml_cuda_mma {
             : "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[3]));
 #endif // __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
 #elif defined(AMD_WMMA_AVAILABLE)
+#if defined(RDNA4)
         using halfx8_t = __attribute__((ext_vector_type(8))) _Float16;
         using floatx8_t = __attribute__((ext_vector_type(8))) float;
         floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
         const halfx8_t& a_frag = reinterpret_cast<const halfx8_t&>(A.x[0]);
         const halfx8_t& b_frag = reinterpret_cast<const halfx8_t&>(B.x[0]);
         acc_frag = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(a_frag, b_frag, acc_frag);
+#endif // RDNA4
 #else
         GGML_UNUSED_VARS(D, A, B);
         NO_DEVICE_CODE;
@@ -873,12 +899,14 @@ namespace ggml_cuda_mma {
     static __device__ __forceinline__ void mma(
             tile<16, 16, float> & D, const tile<16, 8, nv_bfloat162> & A, const tile<16, 8, nv_bfloat162> & B) {
 #if defined(AMD_WMMA_AVAILABLE)
+#if defined(RDNA4)
         using bf16x8_t = __attribute__((ext_vector_type(8))) __bf16;
         using floatx8_t = __attribute__((ext_vector_type(8))) float;
         floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
         const bf16x8_t& a_frag = reinterpret_cast<const bf16x8_t&>(A.x[0]);
         const bf16x8_t& b_frag = reinterpret_cast<const bf16x8_t&>(B.x[0]);
         acc_frag = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(a_frag, b_frag, acc_frag);
+#endif // RDNA4
 #else
         GGML_UNUSED_VARS(D, A, B);
         NO_DEVICE_CODE;
@@ -907,14 +935,14 @@ namespace ggml_cuda_mma {
 #endif // defined(CDNA3)
 
 #elif defined(AMD_WMMA_AVAILABLE)
-        using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int;
-        int32x2_t * a_vec = (int32x2_t *) A.x;
-        int32x2_t * b_vec = (int32x2_t *) B.x;
 
         using int32x8_t = __attribute__((__vector_size__(8 * sizeof(int)))) int;
         int32x8_t * acc = (int32x8_t *) D.x;
 
 #if defined(RDNA4)
+        using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int;
+        int32x2_t * a_vec = (int32x2_t *) A.x;
+        int32x2_t * b_vec = (int32x2_t *) B.x;
 
         acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(
             true,
@@ -933,7 +961,30 @@ namespace ggml_cuda_mma {
             acc[0],
             true
         );
-#endif // defined(RDNA4)
+
+#elif defined(RDNA3)
+        using int32x4_t = __attribute__((__vector_size__(4 * sizeof(int)))) int;
+        int32x4_t * a_vec = (int32x4_t *) A.x;
+        int32x4_t * b_vec = (int32x4_t *) B.x;
+
+        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(
+            true,
+            a_vec[0],
+            true,
+            b_vec[0],
+            acc[0],
+            true
+        );
+
+        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(
+            true,
+            a_vec[1],
+            true,
+            b_vec[1],
+            acc[0],
+            true
+        );
+#endif // RDNA4
 
 #else
         GGML_UNUSED_VARS(D, A, B);
@@ -1020,27 +1071,40 @@ namespace ggml_cuda_mma {
 static __device__ __forceinline__ void mma(
             tile<16, 16, int> & D, const tile<16, 4, int> & A, const tile<16, 4, int> & B) {
 #if defined(AMD_WMMA_AVAILABLE)
-    using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int;
-    int32x2_t * a_vec = (int32x2_t *) A.x;
-    int32x2_t * b_vec = (int32x2_t *) B.x;
-
-    using int32x8_t = __attribute__((__vector_size__(8 * sizeof(int)))) int;
-    int32x8_t * acc = (int32x8_t *) D.x;
-
-    acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(
-        true,
-        a_vec[0],
-        true,
-        b_vec[0],
-        acc[0],
-        false
-    );
+        using int32x8_t = __attribute__((__vector_size__(8 * sizeof(int)))) int;
+        int32x8_t * acc = (int32x8_t *) D.x;
+#if defined(RDNA4)
+        using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int;
+        int32x2_t * a_vec = (int32x2_t *) A.x;
+        int32x2_t * b_vec = (int32x2_t *) B.x;
+
+        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(
+            true,
+            a_vec[0],
+            true,
+            b_vec[0],
+            acc[0],
+            false
+        );
+#elif defined(RDNA3)
+        using int32x4_t = __attribute__((__vector_size__(4 * sizeof(int)))) int;
+        int32x4_t * a_vec = (int32x4_t *) A.x;
+        int32x4_t * b_vec = (int32x4_t *) B.x;
+
+        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(
+            true,
+            a_vec[0],
+            true,
+            b_vec[0],
+            acc[0],
+            false
+        );
+#endif // RDNA4
 #else
         GGML_UNUSED(D);
         GGML_UNUSED(A);
         GGML_UNUSED(B);
         NO_DEVICE_CODE;
-#endif
+#endif // AMD_WMMA_AVAILABLE
     }
 }
-

filename: ggml/src/ggml-cuda/mmq.cu
patch: @@ -307,10 +307,9 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) {
     }
 
     if (amd_wmma_available(cc)) {
-        if (GGML_CUDA_CC_IS_RDNA4(cc)) {
-            return true;
-        }
+        return true;
     }
 
-    return (!GGML_CUDA_CC_IS_RDNA3(cc) && !GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
+    return (!GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE;
+
 }

filename: ggml/src/ggml-cuda/mmq.cuh
patch: @@ -1542,8 +1542,10 @@ static __device__ __forceinline__ void vec_dot_q2_K_q8_1_mma(
             tile_C Cm;
             if (k01 >= MMQ_TILE_NE_K * 3/4) {
                 tile_A A1;
-                A1.x[0] = 0x01010101;
-                A1.x[1] = 0x01010101;
+#pragma unroll
+                for (int l = 0; l < tile_A::ne; ++l) {
+                    A1.x[l] = 0x01010101;
+                }
                 mma(Cm, A1, B);
             }