randomllama commited on
Commit
269cd26
·
verified ·
1 Parent(s): 00d52e8

sync: latest measured numbers

Browse files
Files changed (1) hide show
  1. LADDER.md +11 -6
LADDER.md CHANGED
@@ -15,9 +15,14 @@ code-1/prose-1 prompts as in RESULTS.md. Keep rule: beats incumbent code median
15
 
16
  Round-1 net: +3% code. The big lever (fp8 target KV) is blocked, not disproven.
17
 
18
- ## Round 2 (in progress)
19
- - **L1 per-shape tilelang tiles**: round-1 image shrank ALL DSA kernels to the GB10 tile;
20
- only the 8-token verify shape needed it. Reverting decode-path kernels to stock
21
- (block_I=64/num_stages=2/threads=256), verify-shape kernel stays small.
22
- - L4 autopsy with log capture before teardown.
23
- - Draft-token sweep 8/6/4.
 
 
 
 
 
 
15
 
16
  Round-1 net: +3% code. The big lever (fp8 target KV) is blocked, not disproven.
17
 
18
+ ## Round 2
19
+ | rung | change | code | prose | verdict |
20
+ |---|---|---|---|---|
21
+ | L1v2 | decode+v1 kernels stock tiles | — | — | BOOT FAILED, same 169,984 B smem error — `sparse_attention_fwd_kernel_v1` is in the verify path and must stay small |
22
+ | L1v3 | only `sparse_mla_fwd_decode_partial` stock | 28.5 (25.7–33.2) | 20.2 | kept (wash) — **finding: the GB10 small tile is ~free; tile geometry is not the bottleneck** |
23
+
24
+ Kernel tile map for sm_121 (measured): `qo_len` multi-token kernel → small tile REQUIRED;
25
+ `sparse_attention_fwd_kernel_v1` → small tile REQUIRED (verify path); `sparse_mla_fwd_decode_partial`
26
+ → either (no measurable difference).
27
+
28
+ In progress: draft-token sweep (6 vs 8), L4 fp8-target-KV autopsy with log capture.