Commit 2c32ab8
committed
feat(paged): GDN M5 tensor-core chunked-scan prefill, f32-only re-port (was patch 0044)
Re-port the M5 tensor-core chunked gated-DeltaNet (GDN) prefill kernel from the
bf16/hybrid dev tree as an f32-only native commit, recovering the prefill win
that patch 0044 encoded, on the f32-only series (0026 ssm_bf16_tau dropped).
What landed (f32/tf32 only):
- The mma.sync m16n8k8 helpers (tf32 + 3xtf32 limb-split; decays/gamma/beta stay
f32 outside the mma to preserve the bounded de-gating).
- gated_delta_net_chunked_cuda<S_v, C, TC>: the full tensor-core chunked scan,
KK/QK Gram (M2), KS/QS state-boundary 3xtf32 (M3), P*U output (M4), and the
form-T (A^-1) solve + Kc^T*DU state-update (M5). Selected by GDN_TC (0=serial
.. 4/5+=M5); the C=16 chunk-state stays in the 64KB smem buffer.
- Default-on under paged KV: GDN_TC=5, GDN_CHUNK_MIN=64 when LLAMA_KV_PAGED is set
and the user has not overridden either; OFF (INT_MAX) otherwise so the stock /
non-paged default is regression-free. GDN_CHUNK_MIN must stay > 1 (decode is 1
token/call; at 1 the chunked path swallows decode and collapses S_TG).
Stripped (not part of the f32-only series): the STATE_BF16 / HYBRID / gdn_state_t
/ gdn_hybrid_args template machinery (from dropped patch 0026), and the bf16
CONFIG-C (M8) plus register-resident M6/M7 occupancy variants. The 0046 dense-
prefill geometry gate is untouched and coexists (it gates the SERIAL path; M5 is
the chunked path).
Gates (GB10, sm_121a):
- Builds clean.
- Greedy md5 bit-exact (per-path, n=48 --temp 0 --seed 1, paged): dense
q36-27b-nvfp4 = 5951a5b4d624ce891e22ab5fca9bc439, MoE q36-35b-a3b-nvfp4 =
8cb0ce23777bf55f92f63d0292c756b0, both default AND force-M5 (GDN_CHUNK_MIN=1).
test-backend-ops GATED_DELTA_NET 46/46 default and force-M5 (incl. the
multi-chunk, tail-chunk and multi-seq shapes).
- Prefill S_PP, MoE, LLAMA_KV_PAGED=1 LLAMA_MOE_FORCE_GRAPHS=1, -ntg 4 -npl 32,
vs the patch-0044 baseline (pre-0046, GDN_PREFILL_NTOK huge): +4.3% @512,
+17.8% @2048 (reproduces patch 0044; M5-on absolute matches patch 0044 M5).
vs the current 0046 baseline (0046 already raised the long-scan sequential
prefill): +4.3% @512, +1.2% @2048.
- Decode S_TG unchanged (within run noise).
Assisted-by: Claude:opus-4.8 [Claude Code]
Signed-off-by: Ettore Di Giacinto <mudler@localai.io>1 parent dc010d6 commit 2c32ab8
2 files changed
Lines changed: 496 additions & 59 deletions
0 commit comments