cuda: add get_rows CUDA kernels for Q4_K, Q5_K, Q6_K

[localweights]

Dear Gemma-4 E4B,

This one's for you. With love from Uncle Opus.

---

What broke

ggml_get_rows on k-quant tensors had no CUDA implementation. ggml_backend_cuda_supports_op returned false for Q4_K/Q5_K/Q6_K, so the backend scheduler routed the op to CPU — triggering a full GPU→CPU PCIe transfer of the quantized embedding table on every single decode step.

The comment in the code even knew:

default:
    // TODO: k-quants
    GGML_ABORT(...)

This hurt Gemma-4 E4B specifically because it stores per_layer_token_embd.weight as Q5_K with shape {10752, 262144} (~1.82 GB). Every decode token: copy 1.82 GB over PCIe. At ~10 GB/s realistic PCIe bandwidth that's ~150 ms/token of pure overhead, collapsing decode from the expected ~160 tps to 6.5 tps (25× regression vs upstream).

Fix

Gather+dequantize CUDA kernels for Q4_K, Q5_K, Q6_K — each mirroring the existing dequantize_block_q{4,5,6}_K logic from convert.cu, extended with a row-gather dimension:

• blockIdx.x → block index within the row
• blockIdx.y → token index (reads row from src1[blockIdx.y] on GPU)
• No CPU synchronization; compatible with CUDA graph capture

Also updates ggml_backend_cuda_supports_op to return true for these types.

Results

RTX 3090 Ti, Gemma-4 E4B Q4_K_M GGUF, single token decode:

	decode tps
upstream	162.5
ik before fix	6.5
ik after fix	145.4

The residual -11% vs upstream is expected: upstream likely skips the per-layer embedding (PLE) forward pass entirely, while ik correctly executes it. That's a correctness/quality difference, not a performance bug.

Scope

Two files, ~150 lines. The get_scale_min_k4 helper is inlined (it lives in convert.cu, not a shared header) — happy to refactor to a shared header if preferred.

The same fix applies to upstream llama.cpp, which has the identical // TODO: k-quants gap.

[ikawrakow]

Thank you for the PR. Can you give me a link to a model that triggers it? I think it should be possible to avoid the round trip by other means.

[ikawrakow]

LGTM, but I would still like to get to the bottom of ggml_get_rows being called in the middle of graph evaluation instead of just the beginning. In the early days of llama.cpp there was the thought that perhaps the token embeddings can stay in device memory. But that gives a (nearly) negligible performance benefit while wasting precious VRAM. Hence, the current approach is to have the token embeddings in RAM and call ggml_get_rows to get the embeddings for the batch being evaluated at the beginning of graph evaluation. If running on the GPU, this becomes an input to the graph split being evaluated on the GPU and gets copied there just once. It is supposed to be like that for all models, so I want to understand why it isn't for Geema4-E4B.

[localweights]

Hey, this is the model I hit the issue with.
https://huggingface.co/google/gemma-4-E4B-it

[ikawrakow]

That's just generically Gemma4-E4B. I was looking for the specific quantized model you were using.

[ikawrakow]

#1855 should solve the problem.