You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
Follow-up to #387, which took Carnice (Qwen3.6-35B-A3B-MTP, CPU-MoE) prefill from ~130 → 185.9 t/s (+46%) via the GPU op-offload. This issue records the measured ceiling analysis and the remaining levers to approach llama.cpp's 676 t/s pp512, so future attempts don't re-derive it or retry dead ends.
Baseline (RTX 4070 Ti + Ryzen 9 7900X, same hardware as llama.cpp)
llama-bench -m Carnice.gguf -ngl 99 -ncmoe 99 -fa on:
-nopo 1 (MoE matmul forced on CPU): 282 t/s → so ~2.4× of llama.cpp's lead is GPU op-offload, ~2.2× is a faster CPU kernel.
After #387: op-offload (SHARPI_MOE_GPU_PREFILL=1) = 185.9 t/s (vs our CPU path 127.6). Still ~3.6× under 676.
Why 676 needs a per-layer rewrite, not incremental levers
The op-offload is near its per-layer-serial floor: the residual stream is sequential across the 40 layers (layer L's combine must finish before L+1's trunk reads it). Per-token breakdown @ ~245 tokens (~1700ms total, warm/steady-state):
phase
ms
where
trunk (attn/GDN)
336
GPU, batched
router
376
HOST: matvec + softmax + top-k + shexp-gate, per token per layer
routed-MoE
~720
GPU op-offload; ~600 is the serial floor
combine
~15
host
llama.cpp does ~1.5 ms/tok; we do ~7 ms/tok. Closing it needs the per-layer compute ~3.6× faster.
Remaining levers (ranked)
Router-on-GPU (biggest single, 376ms ≈ 22%). Move the router matvec (norm[N×embDim] × routerW[numExperts×embDim]) + softmax + top-k + shexp-gate onto the GPU. Hard part: the expert bucketing (expStart/expTokI/expTokK CSR) feeds the op-offload's per-expert grouping — it must become GPU-side (or the selection round-trips to host). PrefillBatchedCpuMoe router loop ~CudaHybridGdnForwardPass.cs:2425-2442.
De-serialize the per-layer residual pipeline. The L→L+1 residual dependency serializes the GPU work; only the weight DMA prefetch overlaps. Investigate whether trunk(L+1) can start before MoE(L)'s combine (it can't read the residual, but its norm/QKV could prefetch), or a software-pipelined multi-layer schedule.
Grouped GEMM for the routed experts. Today it's 256 separate per-expert GEMM-N launches/layer (gemm ~256ms). A single grouped/segmented GEMM over all experts (MegaBlocks-style) amortizes launch + improves SM occupancy.
Follow-up to #387, which took Carnice (Qwen3.6-35B-A3B-MTP, CPU-MoE) prefill from ~130 → 185.9 t/s (+46%) via the GPU op-offload. This issue records the measured ceiling analysis and the remaining levers to approach llama.cpp's 676 t/s pp512, so future attempts don't re-derive it or retry dead ends.
Baseline (RTX 4070 Ti + Ryzen 9 7900X, same hardware as llama.cpp)
llama-bench -m Carnice.gguf -ngl 99 -ncmoe 99 -fa on:-nopo 1(MoE matmul forced on CPU): 282 t/s → so ~2.4× of llama.cpp's lead is GPU op-offload, ~2.2× is a faster CPU kernel.After #387: op-offload (
SHARPI_MOE_GPU_PREFILL=1) = 185.9 t/s (vs our CPU path 127.6). Still ~3.6× under 676.Why 676 needs a per-layer rewrite, not incremental levers
The op-offload is near its per-layer-serial floor: the residual stream is sequential across the 40 layers (layer L's combine must finish before L+1's trunk reads it). Per-token breakdown @ ~245 tokens (~1700ms total, warm/steady-state):
llama.cpp does ~1.5 ms/tok; we do ~7 ms/tok. Closing it needs the per-layer compute ~3.6× faster.
Remaining levers (ranked)
norm[N×embDim] × routerW[numExperts×embDim]) + softmax + top-k + shexp-gate onto the GPU. Hard part: the expert bucketing (expStart/expTokI/expTokKCSR) feeds the op-offload's per-expert grouping — it must become GPU-side (or the selection round-trips to host).PrefillBatchedCpuMoerouter loop ~CudaHybridGdnForwardPass.cs:2425-2442.ggml-cpu-zen4is ~2.2× our int8 dots there. perf(cuda): GPU op-offload of MoE prefill (default-on, +46%) — close the Carnice gap to llama.cpp #387's VNNI helper is Q3_K-only and argmax-stable; extending VNNI coverage / a repacked tiled CPU GEMM is the decode lever.Dead ends — do NOT retry (measured, regressed)
(routed+shared*scale)+residto GPU, skip the per-layer routed download): regressed (185.9→139.8). The download was NOT the cost — the per-layer residual-serialization sync is inherent; removing the download just moved it. (perf(cuda): GPU op-offload of MoE prefill (default-on, +46%) — close the Carnice gap to llama.cpp #387 reverted it.)UploadBackgroundRawInto: regressed (+560ms) — that path always stages mmap→pinned synchronously. Use acudaMallocHostbuffer + directcudaMemcpyAsyncinstead (what perf(cuda): GPU op-offload of MoE prefill (default-on, +46%) — close the Carnice gap to llama.cpp #387 does).cudaHostRegisteron the file-backed expert mmap: gives async overlap but only ~13 GB/s (not full PCIe). UsecudaMallocHost(copy once at load) → ~26 GB/s. (perf(cuda): GPU op-offload of MoE prefill (default-on, +46%) — close the Carnice gap to llama.cpp #387.)Measurement discipline (important — easy to get wrong)
cudaMallocHost+ copy + GPU-scratch alloc) was moved to load in perf(cuda): GPU op-offload of MoE prefill (default-on, +46%) — close the Carnice gap to llama.cpp #387 (eagerEnsureGpuOffloadScratchin the ctor) so a single-turn bench reflects steady-state. Before that, cold one-shot was setup-polluted (79–98 t/s vs steady 186).[gpu-offload]per-layer buckets exclude alloc/JIT,routedMoE(in[batched-prefill]) includes them on cold runs.ForceEagerJitentry compiles all), so JIT is at load.Key references
CudaHybridGdnForwardPass.BatchedRoutedExpertsGpuOffload,PrefillBatchedCpuMoe.llm_matvec_q3k_gemm_n(CudaTextKernels.cs), dispatchCudaBackend.MatMulBatched.llm_gather_rows/llm_scatter_rows.EnsureExpertWeightsPinned,UploadRawIntoAsyncDirect,AllocatePinnedHost.SimdKernels.DotQ3K_Q8KS/DotQ4K_Q8KS, nativenative/cpu_vnni/q8k_vnni.c,Q8VnniInterop.🤖 Generated with Claude Code