ggml-cuda : convert contiguous tensors four elements at a time (llama/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

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);
 }