Commit f4e276a20 for llama.cpp

commit f4e276a2066a40cd200db17c1131826d6c0c7a94
Author: Piotr Wilkin (ilintar) <piotr.wilkin@syndatis.com>
Date:   Mon Sep 21 18:00:51 2026 +0200

    ggml-cuda : convert contiguous tensors four elements at a time (#29155)

    convert_unary handles the contiguous case through the general strided kernel,
    one element per thread: each lane reads 4 bytes and writes 2. Converting the
    activations for a bf16 matrix multiplication that way moves 126 MB in 1021 us
    on gfx1151, about 65% of what the memory system can do.

    Give the contiguous path its own kernel that takes four elements per thread
    through a vector type, so a warp loads 512 bytes at a time instead of 128. It
    is used only when the element count is a multiple of four and both pointers
    carry the alignment the vector type needs, and falls back to the strided
    kernel otherwise.

    Model level, Qwen3.8-Next-Flash IQ3_XXS on gfx1151, llama-bench -ub 2048 -r 6,
    mean of the last 3 reps, ABBA counterbalanced:

        pp2048   688.0 680.0  ->  694.3 691.1   +1.26%
        tg128     24.8  24.8  ->   24.8  24.8   +0.14%

    Every conversion in a prefill takes the new kernel (kernel trace: 1146
    convert_unary_cont_vec4, no convert_unary). Output is bit identical; MUL_MAT,
    MUL_MAT_ID, CPY, CONT, GET_ROWS and SET_ROWS pass.

    Assisted-by: Claude Opus 5

diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu
index 360c614a4..0619f4760 100644
--- a/ggml/src/ggml-cuda/convert.cu
+++ b/ggml/src/ggml-cuda/convert.cu
@@ -439,6 +439,29 @@ static __global__ void convert_unary(
     }
 }

+template <typename T> struct alignas(sizeof(T)*4) cvt_vec4 { T v[4]; };
+
+// four elements per thread, so a warp moves 512B (RDNA) / 1k (CDNA) per load
+template <typename src_t, typename dst_t>
+static __global__ void convert_unary_cont_vec4(
+        const void * __restrict__ vx, dst_t * __restrict__ y, const int64_t k4) {
+    const int64_t i = (int64_t)blockDim.x*blockIdx.x + threadIdx.x;
+
+    if (i >= k4) {
+        return;
+    }
+
+    const cvt_vec4<src_t> xv = ((const cvt_vec4<src_t> *) vx)[i];
+
+    cvt_vec4<dst_t> yv;
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        yv.v[j] = ggml_cuda_cast<dst_t>(xv.v[j]);
+    }
+
+    ((cvt_vec4<dst_t> *) y)[i] = yv;
+}
+
 template <typename src_t, typename dst_t>
 static void convert_unary_cuda(const void * vx, dst_t * y,
         const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
@@ -452,6 +475,15 @@ static void convert_unary_cuda(const void * vx, dst_t * y,

 template <typename src_t, typename dst_t>
 static void convert_unary_cont_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
+    if (k % 4 == 0 &&
+        (uintptr_t) vx % alignof(cvt_vec4<src_t>) == 0 &&
+        (uintptr_t) y  % alignof(cvt_vec4<dst_t>) == 0) {
+        const int64_t k4 = k/4;
+        const int64_t num_blocks = (k4 + CUDA_DEQUANTIZE_BLOCK_SIZE - 1) / CUDA_DEQUANTIZE_BLOCK_SIZE;
+        convert_unary_cont_vec4<src_t, dst_t><<<num_blocks, CUDA_DEQUANTIZE_BLOCK_SIZE, 0, stream>>>(vx, y, k4);
+        return;
+    }
+
     convert_unary_cuda<src_t>(vx, y, k, 1, 1, 1, k, k, k, stream);
 }