perf(gemma4): dual-GPU FP8 resident without host OOM + peer mix - #154

Merged

localai-bot merged 54 commits into

mudler:mainfrom

bakon11:feat/gemma4-moe-resident-speed

Aug 8, 2026

Commits on Aug 8, 2026

perf(gemma4): pack resident MoE on GPU0 first then spill GPU1
Don Mirror
committed
597ee89
fix(gemma4): stream FP8 resident upload — no host BF16 cache OOM
Don Mirror
committed
357583a
fix(gemma4): full dual-GPU FP8 resident without host/GPU OOM
Don Mirror
committed
84ee0bf
perf(gemma4): peer-copy off-device resident experts (stay on GPU)
Don Mirror
committed
ce8bfb1
perf(gemma4): device MulScalar+Add for MoE expert mix (stay on GPU)
Don Mirror
committed
4fe4ed7
perf(gemma4): reuse MoE top-k scratch buffers
Don Mirror
committed
fd894a6
perf(gemma4): allocate peer MoE scratch only when needed
Don Mirror
committed
ac1d669
feat(cli): --repeat N for load-once warm tok/s benches
Don Mirror
committed
b4a994b
perf(gemma4): parallel FP8 expert dequant (cold path)
Don Mirror
committed
18d74bf
perf(gemma4): pin host FP8 expert BF16 caches + device-upload diag
Don Mirror
committed
64b3876
perf(gemma4): reuse ExpertGeGLU scratch across top-k
Don Mirror
committed
876eecb
feat(rocm): strided/pointer batched MatmulBT for MoE (opt-in)
Don Mirror
committed
38028ea
perf(gemma4): fused Gelu across top-k in batched MoE (opt-in)
Don Mirror
committed
40166fb
perf(gemma4): hoist MoE scratch outside token loop
Don Mirror
committed
c80838a
feat(rocm): device MoeRouterTopK + Gemma4 uses it
Don Mirror
committed
959d7f6
perf(gemma4): first-expert writes ysum directly (skip Zero+Add)
Don Mirror
committed
8a5464e
perf(gemma4): separate GeLU*up + TLS MoE scratch across layers
Don Mirror
committed
e3bdce0
perf(gemma4): fuse expert mix into down-proj GEMM (alpha/beta)
Don Mirror
committed
f2cc17c
perf(gemma4): PLE GeGLU via GeluMulSeparate (no 2*ple pack)
Don Mirror
committed
5d239fe
perf(gemma4): TLS reuse of per-layer ForwardBody scratch buffers
Don Mirror
committed
ad09bb7
perf(gemma4): contiguous PLE-by-layer layout [L,T,ple]
Don Mirror
committed
dd2463a
perf(gemma4): TLS dense MLP gate_up temps + CLI fflush multi-run stats
Don Mirror
committed
ae2fd08
perf(gemma4): async expert H2D + optional VRAM LRU (no per-expert sync)
Don Mirror
committed
25c9216
perf(gemma4): fuse RmsNorm+residual residual joins (ROCm)
Don Mirror
committed
4895378
perf(gemma4): fuse expert gate+up into one MatmulBT (3→2 GEMMs/expert)
Don Mirror
committed
7bd4b70
perf(gemma4): prefetch top-k expert H2D + fused gate_up in batch path
Don Mirror
committed
5fa5f20
chore(rocm): GemmCompute selector (default 32F; 16bf/16f unsupported on gfx1201 BT)
Don Mirror
committed
2de2457
perf(gemma4): DualRmsNormPlusRes fuse MoE post-FF (3 norms+add+res)
Don Mirror
committed
31eeb00
perf(rocm): thread-local hipBLAS handle (no per-GEMM mutex/SetStream)
Don Mirror
committed
86be92c
fix(rocm): restore mutex/unordered_map includes for batched GEMM tables
Don Mirror
committed
4e8ada6
perf(gemma4): TLS QKV/attn temps in Gemma4AttnBlock
Don Mirror
committed
ed25009
perf(gemma4): TLS router RMS ones-weight + VT_GEMMA4_PROFILE breakdown
Don Mirror
committed
cddf4fe
feat(rocm): opt-in M=1 BF16 GEMV path (default off; hipblas faster)
Don Mirror
committed
43da184
perf(rocm): LDS-cached M=1 BF16 GEMV (still opt-in; hipblas faster)
Don Mirror
committed
381030b
feat(gemma4): fused top-k ExpertGeGLU HIP path (opt-in, correct, slow)
Don Mirror
committed
61f3502
docs: PR draft — Vulkan Q8 ~98 t/s bar vs ROCm FP8 ~34
Don Mirror
committed
ba36e9e
feat(rocm): hipBLASLt MatmulBT path (VT_ROCM_HIPBLASLT=1; default off)
Don Mirror
committed
ab7b124
fix(rocm): hipBLASLt MatmulBT via COL layout (matches GemmEx; no hang)
Don Mirror
committed
e383804
chore(gemma4): disable broken expert batch paths (gather/strided wrong; ptr slow)
Don Mirror
committed
cd2e032
fix(gemma4): resident mode — no FP8 device cache when packs exist; 12GiB GPU0 headroom
Don Mirror
committed
d97290a
feat(gemma4): native FP8 expert path (VT_GEMMA4_FP8_NATIVE=1; default off)
Don Mirror
committed
c6c8531
perf(gemma4): layer profile attn/mlp/moe under VT_GEMMA4_PROFILE
Don Mirror
committed
a897f60
perf(gemma4): top-k fused GeluAndMul (one gelu for all experts)
Don Mirror
committed
f119e9c
perf(gemma4): fused top-k path uses MatmulBTAlphaBetaRocm directly
Don Mirror
committed
dd3f701
docs: gfx1201 FP8 Lt matrix — W8A8 works but slower than BF16 M=1
Don Mirror
committed
1d1e93b
feat(gemma4): custom RDNA4 BF16 Expert GeGLU kernels (VT_GEMMA4_CUSTOM_EXPERT)
Don Mirror
committed
f4436f8
feat(server): load chat_template.jinja + verbose chat request logs
Don Mirror
committed
f42f4ff
feat(server): chat-dbg stage heartbeats while processing prompts
Don Mirror
committed
f162d93
feat(server): vLLM-style request logger + metrics for Hermes serve
Don Mirror
committed
5f70f66
fix(server): reject huge Hermes SOUL prompts; clamp max_tokens
Don Mirror
committed
98b2bbd
feat(server): Hermes long-ctx defaults — 200k prompt cap, APC note
Don Mirror
committed
0483ef9
feat(sched): log chunked prefill progress (INFO prefill computed=N/M %)
Don Mirror
committed
0bae7cb
fix(engine): generate waits on own request_id; chat enable_thinking flag
Don Mirror
committed
3cbecb5
fix(gemma4): portable vt::fused_ops seam — no vt::rocm in models/
Don Mirror
committed
e086435