P3b-2 for the Blackwell W4A16 Marlin GEMM. The q4_K dequant wall is partly
cross-N-block-redundant: every N-block re-decodes the same weight strip, so
halving the N-block count (BN 64->128) halves that redundant 6-bit superblock
decode. A BN sweep showed this only pays off when BN is spread across more
warps (16 warps, 8 m16n8 C-tiles/warp) rather than more fragments-per-warp -
the FN=8 / FM=4 variants (16 C-tiles/warp) regressed to ~6.6 TFLOPS on
register pressure. Shipping tile is now WM=4,WN=4,FM=2,FN=4 -> BM=128, BN=128,
16 warps.
Thermally-bracketed cold A/B (q4_K n=512 / q4_0 n=512 via test-backend-ops
perf; pp512/pp2048 via llama-bench Qwen3-32B-Q4_K_M):
BN64/8w (prev): 8.50 / 10.56 TFLOPS, measured 8.45/10.51 again (bracket)
BN128/16w (this): 9.92 / 11.68 TFLOPS, pp512 177.6, pp2048 185.0
-> +17% q4_K, +11% q4_0, +20% pp512 vs the previous commit; +49% pp512 vs
the original block-tiled kernel (119).
Parity gate GGML_CUDA_W4A16=1 test-backend-ops MUL_MAT = 1103/1103, flag set
and unset (byte-identical when unset). Still ~4.7x under MMQ (47 TFLOPS) and
does NOT beat MMQ; BN growth divides the redundant decode but cannot remove
the per-k-step decode itself - the offline weight prepack remains the next
unlock for q4_K. Plan doc P3 table + bottleneck notes updated.
Assisted-by: Claude:opus-4.8 [Claude Code]
Signed-off-by: Ettore Di Giacinto <mudler@localai.io>
P3b for the Blackwell (sm_120/121) W4A16 Marlin GEMM. Two combined changes
over the prior block-tiled kernel, both verified by a thermally-bracketed
cold A/B (committed measured identically before and after):
- Skew-padded shared layout: store the staged weight/activation rows at a
padded stride of 12 bf162 (8 data + 4 pad) and feed the tensor cores with
ldmatrix.x4 (A) / ldmatrix.x2 (B). ldmatrix's per-lane address is
row*stride; the natural stride 8 divides the 32-bank cycle and collides
rows 0,4,8,12 (2-way bank conflict). Skewing to 12 (still 16-byte aligned)
spreads {r*12 mod 32} across 8 distinct bank-quads, so both ldmatrix halves
are conflict-free at only +50% on the ~6 KB staged tile - unlike a 128-byte
-row XOR swizzle, which is conflict-free but needs 16 KB shared and
collapses occupancy on GB10 (measured 2.84 TFLOPS, worse than baseline).
- Larger tile: BM=128, BN=64, 8 warps (WM=4,WN=2,FM=2,FN=4), which cuts the
redundant per-M-block activation re-reads.
Cold A/B (q4_K n=512 / q4_0 n=512 via test-backend-ops perf; pp512/pp2048 via
llama-bench Qwen3-32B-Q4_K_M):
committed: 6.63 / 7.53 TFLOPS, pp512 119
this: 8.52 / 10.49 TFLOPS, pp512 148.5, pp2048 153.9 (+28% / +40% / +25%)
Parity gate GGML_CUDA_W4A16=1 test-backend-ops MUL_MAT = 1103/1103, flag set
and unset (byte-identical when unset). Still ~5.5x under MMQ (47 TFLOPS) and
does NOT beat MMQ yet; the q4_K limiter has now moved from the mma feed to the
per-element 6-bit superblock dequant (q4_0 scales to 15.8 TFLOPS with more
warps while q4_K stays ~8.5), so the offline weight prepack is the next unlock.
Plan doc P3 section updated with the sweep data and the corrected bottleneck.
Assisted-by: Claude:opus-4.8 [Claude Code]
Signed-off-by: Ettore Di Giacinto <mudler@localai.io>
Replace the P2 1-warp-per-16x8 W4A16 kernel with a block-tiled multi-warp
kernel: blockDim=(32, WM*WN) so threadIdx.x is the warp lane (required by
mma.cuh get_i/get_j) and threadIdx.y is the warp index. WM*WN warps compute a
BM(=WM*FM*16) x BN(=WN*FN*8) output tile, each warp owning an FM x FN grid of
m16n8k16 BF16 mma fragments accumulated in F32. The BM x 16 dequantized Q4
weight strip is staged once per k-step in a small (~4 KB) shared buffer and
reused across the block's whole BN span. Shipping config WM=2,WN=2,FM=2,FN=4.
The P2 launch put all threads on threadIdx.x; with >1 warp that drove the mma
tile get_j past the shared bound (out-of-bounds shared read, caught by
compute-sanitizer). The new (32, nwarps) layout matches mmf.cu and fixes it.
Parity gate holds 1103/1103 (test-backend-ops MUL_MAT CUDA0), flag set and
unset (byte-identical when GGML_CUDA_W4A16 is unset; the seam returns false).
Perf (q4_K m=4096 k=14336 n=512): ~2 TFLOPS (P2) -> ~7-9 TFLOPS (thermal
dependent); llama-bench Qwen3-32B-Q4_K_M pp512 31.75 -> ~118-142 t/s. Still
below the MMQ baseline (47 TFLOPS / 718 t/s): a tile sweep stayed flat and
q4_0 vs q4_K differ by only ~12%, so dequant compute is not the limiter - the
shared-load / mma-feed is. A naive double-buffered cp.async pipeline (32 KB
shared) regressed via occupancy collapse and an ldmatrix swap was neutral
(unswizzled layout bank-conflicts), both reverted. The path to >=150 TFLOPS is
the full Marlin machinery (XOR-swizzled shared layout + offline weight reshuffle
+ tuned async pipeline + Stream-K), deferred to P3 step 4. See
W4A16_MARLIN_KERNEL_PLAN.md for the per-step table and dead-end notes.
Assisted-by: Claude:opus-4.8 [Claude Code]
Signed-off-by: Ettore Di Giacinto <mudler@localai.io>
Replace the P1 dispatch-seam TODO in marlin-w4a16.cu with a real W4A16
GEMM for consumer Blackwell (sm_120/121). In-kernel dequant of Q4 weights
to BF16, mma.sync m16n8k16 f32.bf16.bf16.f32 tensor-core multiply against
BF16-converted f32 activations, f32 accumulate and write, reusing ggml's
mma.cuh tile abstractions.
Handles the contiguous 2D GEMM prefill path for Q4_0 and Q4_K (f32
activations, ne2==ne3==1); batched, broadcast, permuted, non-contiguous
and f16-activation cases return false and fall back to MMQ so the gate
stays green. M/N boundaries are zero-padded in-kernel.
Parity gate (GGML_CUDA_W4A16=1 test-backend-ops MUL_MAT on GB10):
1103/1103 passed; default flag-off build stays byte-identical 1103/1103.
Model sanity: Qwen3-32B-Q4_K_M llama-bench pp512 31.75 t/s (slow is
expected for P2 - the naive single-warp kernel is the correctness
checkpoint; P3 adds the cp.async pipeline and weight reshuffle).
Assisted-by: Claude:opus-4.8 [Claude Code]
Signed-off-by: Ettore Di Giacinto <mudler@localai.io>
Same strategy as P2: one fresh Opus-4.8 subagent per phase, each handed a
complete zero-context brief, dispatched sequentially as each predecessor lands
(P3 pipeline needs P2's correct kernel, P4 tune needs P3, P5 enable needs P4).
Shared DGX/harness/commit boilerplate factored into a COMMON section; each phase
brief carries its goal, incremental steps, acceptance gate, and a splice note for
the prior phase's actual deliverable.
Assisted-by: Claude:opus-4.8 [Claude Code]
Signed-off-by: Ettore Di Giacinto <mudler@localai.io>
