cuda: PQ2_0 mat-vec kernel for 3-8 columns on Ada

[sb32445]

Overview

Adds a mat-vec kernel for PQ2_0 weights with 3-8 activation columns on Ada (cc 8.9). The generic mmvq kernel drops to 36-66 % of DRAM bandwidth from 4 columns on. The new kernel feeds the raw 2-bit codes straight into dp4a with an integer correction term (new activation layout GGML_CUDA_Q8_1_PQ2), 4 rows per warp, __launch_bounds__(128, 3).

RTX 4070, m=5120 k=17408 (cold, ncu): 4 columns 105.6 -> 54.0 us, 8 columns 136.8 -> 78.1 us.
Ternary-Bonsai-2-27B-PQ2_0, parallel decode sequences: 4: +17 %, 6: +25 %, 8: +29 %.

Additional information

• Only Ada, PQ2_0, 3-8 columns, no fusion / ids / batch dims. Everything else takes the old path.
• The branch has two commits: the first contains the environment switch GGML_CUDA_PQ2_MULTICOL (=0 switches the kernel off) that was used for the measurements below, the last one removes it. To reproduce a measurement, build the first commit.
• test-backend-ops MUL_MAT 219/219 (CUDA vs CPU, incl. odd row counts). Perplexity with and without the path 6.378 (difference below 0.07 %, error +-0.53).
• Not measured on other GPUs. The condition is cc == GGML_CUDA_CC_ADA_LOVELACE exactly (8.9), the registers (168, 3 blocks per SM) are tuned for that GPU.
• Not bit-identical to the generic kernel (the integer part is exact, but the float accumulation order differs); details below. Within the new kernel the result of a column does not depend on how many columns share the call (3 to 8), which the generic kernel does not guarantee.
• ggml_cuda_q8_1_layout_host got a plain_2d argument; the layout choice for the quantizer and the kernel choice come from the same y_layout, and the PQ2 branch asserts that no ids, batch dims or fusion are involved (fusion with more than one column does not exist for PQ2_0).
• PTQ1_0 is not affected (own kernel).

Test results

• Hardware / software: RTX 4070 12 GB (AD104, cc 8.9, 504 GB/s, 48 MB L2), Linux 6.18, NVIDIA driver 615.71, CUDA 13.4, GCC 16.2; Release build, -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=89.
• Base: speed numbers were measured on prism at 88c4bc60b; the four commits since (SYCL, WebGPU and cuda: fused FWHT quantizer for 64-wide warps (#303)) touch quantize.cu only in the 64-wide-warp FWHT path, not the code of this PR. The branch is rebased on 2459f68b5, builds, and test-backend-ops was repeated on it.
• Model: Ternary-Bonsai-2-27B in PQ2_0 (m=5120 x k=17408 is one FFN shape of it).
• test-backend-ops test -b CUDA0 -o MUL_MAT: 1586/1586 passed on the rebased branch (CUDA0 against CPU; includes the PQ2_0 cases with odd row counts). On the earlier base the PQ2_0 subset alone was 219/219.
• Kernel, cold L2 (ncu), m=5120 k=17408: 4 columns 105.6 -> 54.0 us, 8 columns 136.8 -> 78.1 us; the generic kernel reaches 36-66 % of the DRAM bandwidth from 4 columns on (46 % at 4, 36 % at 8 columns).
• Decode of Ternary-Bonsai-2-27B-PQ2_0 with parallel sequences (llama-batched-bench): 4 sequences +17 %, 6 sequences +25 %, 8 sequences +29 %.
• Perplexity (256 context, 6 chunks, -ub 8 / 5 / 4) before -> after: 6.3784 / 6.3784 / 6.3776 -> 6.3774 / 6.3774 / 6.3820, a difference of at most 0.07 %, far inside the reported error of +-0.53.
• Numerics, greedy decode with 4 parallel sequences (-np 4, 4 different prompts sent at once, 192 tokens, q4_0 K/V, two rounds per setting): with the new kernel both rounds are identical for all 4 prompts. With the generic kernel (GGML_CUDA_PQ2_MULTICOL=0) two identical runs already differ for 3 of 4 prompts (the arithmetic of the generic kernel depends on the number of columns, and the batch composition varies slightly between runs). Old against new: 2 of 4 prompts identical, the others first differ at character 161 of 848 and 407 of 838.
• Not tested: other GPUs, HIP, MoE (ids) models, column counts outside 3 to 8 (they take the old path), a task-level quality metric.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: The patches were developed with Claude Code (Anthropic's coding agent): it wrote the code, the measurement scripts and the first drafts of the commit messages and PR texts. I decided what to work on (which kernels and host paths to optimise, based on profiles of my own decode setup). The measurements and checks listed in the PR texts were run in the Claude Code sessions; I did not re-run them independently. I will maintain the changes. Commits where Claude Code was used carry a Co-Authored-By trailer.