RecursiveSI

/

flint

Like 1.24k

Follow

Recursive Superintelligence, Inc. 652

Numbers from a single RTX 3090 (24 GB) with llama.cpp, for anyone sizing Flint for a 24 GB NVIDIA card. I also found one CUDA kernel that takes about 14% of decode time; a small patch for it is at the end.

Setup: RTX 3090 at its stock 350 W limit, in a USB4 enclosure (this only matters for the expert-offload row). Host: Ryzen AI Max+ 395. llama.cpp build 11954 (f1a7dea99), CUDA 13.1, GGUFs from unsloth/flint-GGUF. A 4 GB TTS model shares the card, so I kept Flint under about 19 GiB. llama-bench with -ngl 999 -fa on; each number is the average of 2-3 runs.

Through llama-server (16K context, f16 KV, with the fix, thinking off) it averaged 156 tok/s over a mix of agentic, chat, code, math, rewrite and summarize prompts.

What made a difference

• -ub 2048 for long prompts: pp4096 went from 3497 to 6066 tok/s (+73%). It needs more VRAM; -ub 4096 ran out of memory here next to the TTS model.
• f16 KV is faster than q8_0: at 16K tokens deep, 146.6 vs 122.4 tok/s. I only use q8_0 to fit 32K context.
• With experts in system RAM and a slow GPU link, add --no-op-offload. Otherwise llama.cpp copies those experts to the GPU for every large batch. Here prompt speed went from 272 to 1121 tok/s (pp512), with the same decode speed and 1 GiB less VRAM. On a full PCIe slot the default may be fine; I haven't tested that.
• CUDA beat Vulkan by 14% on decode before the fix and 31% after it.
• --spec-type ngram-mod gave +1% overall in llama-server. It only helps when text repeats, as in code edits.

The slow kernel

An nsys profile of decode showed 7.11 ms per token. 1.02 ms of that (14.4%) was pool1d_nchw_kernel: 47 launches per token, one per MoE layer, about 22 µs each. It is the max-pool behind the MoE amax rescaling that runs before each expert down-projection. The CUDA kernel uses one thread per output, so each 512-wide row is scanned by a single thread, and a decode step has only 6 such rows per layer.

Metal got a fix for this in llama.cpp PR 30862. CUDA has not: ggml/src/ggml-cuda/pool1d.cu has not changed since August.

The patch below gives each row one warp:

• Kernel time on Flint's decode shape went from 22.0 to 1.4 µs (test-backend-ops perf).
• Decode went from 139.8 to 161.4 tok/s (+15.5%); prompt speed rose 2%.
• Output is unchanged. Wikitext perplexity matches chunk for chunk (24.3096 both), and test-backend-ops -o POOL_1D passes 226/226, including whole-row cases I added.

Setting GGML_CUDA_POOL1D_SLOW=1 switches back to the old kernel for A/B tests. I tested it on a 3090 only. It applies at f1a7dea99.

diff --git a/ggml/src/ggml-cuda/pool1d.cu b/ggml/src/ggml-cuda/pool1d.cu
index ac6fb0cbd..51901e128 100644
--- a/ggml/src/ggml-cuda/pool1d.cu
+++ b/ggml/src/ggml-cuda/pool1d.cu
@@ -49,6 +49,24 @@ static __global__ void pool1d_nchw_kernel(
     o_ptr[cur_ow] = res;
 }
 
+// one warp per row for a max over the whole row (e.g. MoE amax rescaling); max is order-independent
+static __global__ void pool1d_row_max_kernel(const int iw, const int nr, const float * src, float * dst) {
+    const int row  = blockIdx.x * (blockDim.x / WARP_SIZE) + threadIdx.x / WARP_SIZE;
+    const int lane = threadIdx.x % WARP_SIZE;
+    if (row >= nr) {
+        return;
+    }
+    const float * x = src + (int64_t) row * iw;
+    float res = -FLT_MAX;
+    for (int i = lane; i < iw; i += WARP_SIZE) {
+        res = fmaxf(res, __ldg(x + i));
+    }
+    res = warp_reduce_max(res);
+    if (lane == 0) {
+        dst[row] = res;
+    }
+}
+
 static void pool1d_nchw_kernel_f32_f32_cuda(
         const int iw, const int ow,
         const int kw, const int sw, const int pw,
@@ -81,5 +99,13 @@ void ggml_cuda_op_pool1d(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
 
     const int parallel_elements = (int)(nr * OW);
 
+    static const bool lab_slow = getenv("GGML_CUDA_POOL1D_SLOW") != nullptr; // lab A/B switch
+    if (!lab_slow && op == GGML_OP_POOL_MAX && OW == 1 && p0 == 0 && k0 >= IW && ggml_is_contiguous(src0)) {
+        const int rows_per_block = 256 / WARP_SIZE;
+        const int num_blocks = (int)((nr + rows_per_block - 1) / rows_per_block);
+        pool1d_row_max_kernel<<<num_blocks, 256, 0, stream>>>((int) IW, (int) nr, src0_d, dst_d);
+        return;
+    }
+
     pool1d_nchw_kernel_f32_f32_cuda(IW, OW, k0, s0, p0, parallel_elements, src0_d, dst_d, op, stream);
 }

I did the measurements and the patch with an AI coding assistant (Claude Code). Every number above comes from a logged run.