- go →
- auto
- auto
The two axes of hard
Every kernel people speak of with reverence (flash, Splash, ragged paged attention, fused MoE) is hard for one or both of exactly two reasons. Axis one: algorithmic restructuring, turning a multi-pass computation into a streaming one, which is a theorem, not a schedule. Axis two: data-dependent iteration, where the loop structure itself depends on runtime data: a sparsity mask, a page table, routing decisions. Master both axes on their simplest instances and the priesthood kernels stop being magic; that is this stage.
The softmax monoid
The central result of axis one is online softmax. The naive form needs the row max before the exponentials, hence the multi-pass shape and stage 2's spill. The streaming form carries a triple (m, l, acc): running max, running rescaled denominator, running rescaled output, and folds each new block in with an exact rescale, since . The deep fact you prove in LAB·3.1 is that the combine operation on these triples is associative and commutative. That single proof is why blocked, streaming, parallel, and (in stage 4) distributed attention are all simultaneously correct. EX·03 plays the schedule with the state updates written as exact formulas.
One associativity proof makes blocked, streaming, and distributed attention correct at once.
The backward, streamed
A fast forward with a slow backward is useless for training, so the backward gets its own week. Two ideas carry it: the identity collapses the softmax Jacobian into something streamable, and recomputing P blockwise from the saved (m, l) beats storing it, trading cheap FLOPs for expensive bytes. You wire your forward to a streaming backward with custom_vjp and let jax.grad referee every step.
Masks become loop structure
Axis two starts with the observation that a causal mask is not something you add to scores; it is blocks you never visit. In EX·08 the kv loop bound depends on the query block index, so blocks past the diagonal are never loaded at all: the mask became loop structure, and the speedup scales with sparsity. Splash attention is this idea industrialized (arbitrary block masks via scalar prefetch), and ragged paged attention is the same axis with a page table as the data. Read both after you have built the toy; the diffs are the syllabus.
Derive before you read
One rule for the whole stage: derivation before code, your own build before any reference. The blind-build discipline is not ceremony; reading Splash first would rob you of the exact struggle that makes its choices legible.
Instruments
01/06prologue · Q pins resident in VMEM; K0, V0 stream in from HBM
01/05prologue · query block 2 pins resident; only blocks on or before the diagonal will ever load
Guided walks
1# one query block per grid step; the kv axis streams inside the kernel2def flash_kernel(q_ref, k_ref, v_ref, o_ref, *, block_kv):3 q = q_ref[...].astype(jnp.float32)4 n_kv = k_ref.shape[0]56 def step(j, state):7 m, l, acc = state8 kb = jax.lax.dynamic_slice_in_dim(k_ref[...], j * block_kv, block_kv).astype(jnp.float32)9 vb = jax.lax.dynamic_slice_in_dim(v_ref[...], j * block_kv, block_kv).astype(jnp.float32)10 s = q @ kb.T11 m_new = jnp.maximum(m, jnp.max(s, axis=-1))12 alpha = jnp.exp(m - m_new)13 p = jnp.exp(s - m_new[:, None])14 l_new = l * alpha + jnp.sum(p, axis=-1)15 acc_new = acc * alpha[:, None] + p @ vb16 return m_new, l_new, acc_new1718 m0 = jnp.full((q.shape[0],), -jnp.inf, jnp.float32)19 l0 = jnp.zeros((q.shape[0],), jnp.float32)20 acc0 = jnp.zeros((q.shape[0], v_ref.shape[1]), jnp.float32)21 m, l, acc = jax.lax.fori_loop(0, n_kv // block_kv, step, (m0, l0, acc0))22 o_ref[...] = (acc / l[:, None]).astype(o_ref.dtype)2324import functools2526def flash(q, k, v, block_q=128, block_kv=128):27 sq, d = q.shape28 return pl.pallas_call(29 functools.partial(flash_kernel, block_kv=block_kv),30 grid=(sq // block_q,),31 in_specs=[32 pl.BlockSpec((block_q, d), lambda i: (i, 0)),33 pl.BlockSpec(k.shape, lambda i: (0, 0)),34 pl.BlockSpec(v.shape, lambda i: (0, 0)),35 ],36 out_specs=pl.BlockSpec((block_q, d), lambda i: (i, 0)),37 out_shape=jax.ShapeDtypeStruct(q.shape, q.dtype),38 interpret=INTERP,39 )(q, k, v)
01/08flash_kernel takes four refs plus a keyword-only block_kv, the size of each key/value chunk the loop below pulls one at a time. The comment above the signature states the design in one line: one query block lives in VMEM for the whole kernel call, and it's the key/value axis that streams through in pieces. Casting q to float32 here is algorithm, not schedule: the running max-and-sum recurrence that follows needs full precision to stay numerically stable.
1import functools23def causal_flash_kernel(q_ref, k_ref, v_ref, o_ref, *, block_q, block_kv):4 qi = pl.program_id(0)5 q = q_ref[...].astype(jnp.float32)6 row0 = qi * block_q78 def step(j, state):9 m, l, acc = state10 kb = jax.lax.dynamic_slice_in_dim(k_ref[...], j * block_kv, block_kv).astype(jnp.float32)11 vb = jax.lax.dynamic_slice_in_dim(v_ref[...], j * block_kv, block_kv).astype(jnp.float32)12 s = q @ kb.T13 cols = j * block_kv + jax.lax.broadcasted_iota(jnp.int32, s.shape, 1)14 rows = row0 + jax.lax.broadcasted_iota(jnp.int32, s.shape, 0)15 s = jnp.where(cols <= rows, s, -jnp.inf) # edge block: mask inside16 m_new = jnp.maximum(m, jnp.max(s, axis=-1))17 alpha = jnp.exp(m - m_new)18 p = jnp.exp(s - m_new[:, None])19 l_new = l * alpha + jnp.sum(p, axis=-1)20 acc_new = acc * alpha[:, None] + p @ vb21 return m_new, l_new, acc_new2223 # the causal structure IS the loop bound: blocks past the diagonal never run24 n_live = (row0 + block_q + block_kv - 1) // block_kv25 m0 = jnp.full((block_q,), -jnp.inf, jnp.float32)26 l0 = jnp.zeros((block_q,), jnp.float32)27 acc0 = jnp.zeros((block_q, v_ref.shape[1]), jnp.float32)28 m, l, acc = jax.lax.fori_loop(0, n_live, step, (m0, l0, acc0))29 o_ref[...] = (acc / l[:, None]).astype(o_ref.dtype)
01/07causal_flash_kernel takes the same four refs as flash_kernel plus one more static argument: block_q joins block_kv, because this kernel now needs to know both how wide a query tile is and how wide a key/value chunk is. Wiring the arguments through functools.partial and BlockSpecs works exactly as it did for plain flash attention; what changes is what happens inside.
1def _shrink_mask_info(2 *,3 block_mask: np.ndarray,4 data_next: np.ndarray,5 mask_next: np.ndarray | None,6 head_shards: int,7):8 assert block_mask.ndim == 39 assert data_next.ndim == 310 assert mask_next is None or mask_next.ndim == 31112 head_block_mask = block_mask[0]1314 grouped_non_zero_cols = []15 # Group non-zero columns based on which row they belong to.16 for row_index in range(head_block_mask.shape[0]):17 head_block_mask_row = head_block_mask[row_index, :]18 non_zero_cols = np.nonzero(head_block_mask_row)[0]19 grouped_non_zero_cols.append(non_zero_cols)2021 # Pad each row in the non-zero indices to match the width of the longest22 # row. This avoids having jagged rows.23 max_non_zero_cols = max(len(x) for x in grouped_non_zero_cols)24 padded_non_zero_cols_list = []25 padding = -126 for row in grouped_non_zero_cols:27 padded_non_zero_cols_list.append(28 np.pad(29 row,30 pad_width=(0, max_non_zero_cols - row.shape[0]),31 constant_values=padding,32 )33 )3435 padded_non_zero_cols = np.stack(padded_non_zero_cols_list, axis=0)3637 assert padded_non_zero_cols.shape[0] == block_mask.shape[1], (38 padded_non_zero_cols.shape,39 block_mask.shape,40 )4142 # For each row of array, select the columns indices in padded_non_zero_cols,43 # ignore padding.44 def select_cols(array):45 assert array.ndim == 246 assert padded_non_zero_cols.ndim == 247 assert array.shape[0] == padded_non_zero_cols.shape[0]48 assert array.shape[1] >= padded_non_zero_cols.shape[1]49 selected_rows = []50 for row in range(array.shape[0]):51 col = padded_non_zero_cols[row]52 selected = array[row][col]53 selected = np.where(col != padding, selected, 0)54 selected_rows.append(selected)5556 return np.stack(selected_rows, axis=0)5758 return _slice_mask_info(59 block_mask=block_mask,60 data_next=data_next,61 mask_next=mask_next,62 head_shards=head_shards,63 slice_function=select_cols,64 )
01/08_shrink_mask_info takes three arrays your causal kernel never had to build by hand: block_mask, data_next, and mask_next. block_mask is a mask already reduced to block granularity, one entry per (head, q_block, kv_block) holding 0 if that block is all zeros, 1 if it mixes zeros and ones, or 2 if it's entirely live. data_next and mask_next are pointer arrays, computed earlier in this same file, that tell the kernel which kv block and which partial-mask block to fetch next. This function's job is to drop the blocks nobody needs to visit, the same job n_live did in your causal kernel, but for a mask of any shape, not just a triangle.
What gets built
- Online softmax derived on paper from a blank page, associativity proven
- Flash attention forward, written from the derivation before reading any reference
- The backward pass, derived and wired with custom_vjp, differential-tested
- Toy block-sparse and ragged kernels; deep reads of Splash, ragged paged attention, MoE dispatch
Readings
- Splash attention source · only after your blind build passes; diff yours against it and catalog every difference
- Ragged Paged Attention paper · axis two at production scale: the page table becomes the loop
- tpu-inference kernels · the serving kernels in production; week 8 reading
The work, in order
week 6 · flash, derived then built
- On paper, from blank page: two-pass softmax to the `(m, l)` pair to the `(m, l, acc)` triple; prove the combine associative and commutative
- LAB·3.1: check the derivation numerically, including 20 random reduction bracketings
- LAB·3.2: the blind build; no references until check passes, then diff against Splash and Tokamax
week 7 · the backward
- Derive dQ, dK, dV on paper; find the rowsum identity yourself before confirming it
- LAB·3.3: wire custom_vjp, differential-test values and grads against jax.grad
- Upgrade the whole-array backward to a blocked streaming pass; corner cases: fully-masked row, length-1 sequence
week 8 · data shapes the loop
- LAB·3.4: causal flash where the loop bound is the mask; measure the sparsity scaling
- Extend to an arbitrary block mask passed as data; you have now rebuilt Splash's core idea
- Read the RPA paper and kernel: watch the page table and lengths become loop structure via scalar prefetch
Labs
Deriving online softmax
Flash attention, blind build
The backward pass
Skipping blocks: masks as loop structure
From the notebook
# the state monoid: combine two partial (m, l, acc) summaries exactly
def combine(a, b):
m1, l1, acc1 = a
m2, l2, acc2 = b
m = jnp.maximum(m1, m2)
a1, a2 = jnp.exp(m1 - m), jnp.exp(m2 - m)
return (m, l1 * a1 + l2 * a2, acc1 * a1[..., None] + acc2 * a2[..., None])
def summarize(s_block, v_block):
m = jnp.max(s_block, axis=-1)
p = jnp.exp(s_block - m[..., None])
return (m, jnp.sum(p, axis=-1), p @ v_block)
def streaming_attention(q, k, v, block=64):
s = q @ k.T
n = s.shape[-1]
state = summarize(s[:, :block], v[:block])
for j in range(block, n, block):
state = combine(state, summarize(s[:, j:j+block], v[j:j+block]))
m, l, acc = state
return acc / l[..., None]import random
def reduce_bracketed(states, order):
states = list(states)
for idx in order:
merged = combine(states[idx], states[idx + 1])
states[idx:idx + 2] = [merged]
return states[0]
s = q @ k.T
blocks = [summarize(s[:, j:j+64], v[j:j+64]) for j in range(0, 512, 64)]
base = reduce_bracketed(blocks, [0] * (len(blocks) - 1)) # left fold
for trial in range(20):
order = [random.randrange(n) for n in range(len(blocks) - 1, 0, -1)]
m, l, acc = reduce_bracketed(blocks, order)
check(f"bracketing {trial}", acc / l[..., None], base[2] / base[1][..., None], tol=1e-5)
print("combine is associative under 20 random bracketings")# one query block per grid step; the kv axis streams inside the kernel
def flash_kernel(q_ref, k_ref, v_ref, o_ref, *, block_kv):
q = q_ref[...].astype(jnp.float32)
n_kv = k_ref.shape[0]
def step(j, state):
m, l, acc = state
kb = jax.lax.dynamic_slice_in_dim(k_ref[...], j * block_kv, block_kv).astype(jnp.float32)
vb = jax.lax.dynamic_slice_in_dim(v_ref[...], j * block_kv, block_kv).astype(jnp.float32)
s = q @ kb.T
m_new = jnp.maximum(m, jnp.max(s, axis=-1))
alpha = jnp.exp(m - m_new)
p = jnp.exp(s - m_new[:, None])
l_new = l * alpha + jnp.sum(p, axis=-1)
acc_new = acc * alpha[:, None] + p @ vb
return m_new, l_new, acc_new
m0 = jnp.full((q.shape[0],), -jnp.inf, jnp.float32)
l0 = jnp.zeros((q.shape[0],), jnp.float32)
acc0 = jnp.zeros((q.shape[0], v_ref.shape[1]), jnp.float32)
m, l, acc = jax.lax.fori_loop(0, n_kv // block_kv, step, (m0, l0, acc0))
o_ref[...] = (acc / l[:, None]).astype(o_ref.dtype)import functools
def flash(q, k, v, block_q=128, block_kv=128):
sq, d = q.shape
return pl.pallas_call(
functools.partial(flash_kernel, block_kv=block_kv),
grid=(sq // block_q,),
in_specs=[
pl.BlockSpec((block_q, d), lambda i: (i, 0)),
pl.BlockSpec(k.shape, lambda i: (0, 0)),
pl.BlockSpec(v.shape, lambda i: (0, 0)),
],
out_specs=pl.BlockSpec((block_q, d), lambda i: (i, 0)),
out_shape=jax.ShapeDtypeStruct(q.shape, q.dtype),
interpret=INTERP,
)(q, k, v)
q = jax.random.normal(jax.random.key(0), (512, 64), jnp.float32)
k = jax.random.normal(jax.random.key(1), (1024, 64), jnp.float32)
v = jax.random.normal(jax.random.key(2), (1024, 64), jnp.float32)
ref = jax.nn.softmax(q @ k.T, axis=-1) @ v
check("flash forward", flash(q, k, v), ref, tol=1e-4)# reference math and the custom_vjp pair; forward saves only (O, m, l) style residuals
def attention_ref(q, k, v):
return jax.nn.softmax(q @ k.T, axis=-1) @ v
@jax.custom_vjp
def attention(q, k, v):
return attention_ref(q, k, v)
def fwd(q, k, v):
s = q @ k.T
m = jnp.max(s, axis=-1, keepdims=True)
p = jnp.exp(s - m)
l = jnp.sum(p, axis=-1, keepdims=True)
o = (p / l) @ v
return o, (q, k, v, o, m, l)
def bwd(res, do):
q, k, v, o, m, l = res
# recompute P blockwise in the real kernel; whole-array here for clarity
p = jnp.exp(q @ k.T - m) / l
dv = p.T @ do
d = jnp.sum(do * o, axis=-1, keepdims=True) # the identity that kills the Jacobian
dp = do @ v.T
ds = p * (dp - d)
return ds @ k, ds.T @ q, dv
attention.defvjp(fwd, bwd)q = jax.random.normal(jax.random.key(0), (128, 64))
k = jax.random.normal(jax.random.key(1), (256, 64))
v = jax.random.normal(jax.random.key(2), (256, 64))
def loss_custom(q, k, v): return jnp.sum(jnp.tanh(attention(q, k, v)))
def loss_ref(q, k, v): return jnp.sum(jnp.tanh(attention_ref(q, k, v)))
g1 = jax.grad(loss_custom, argnums=(0, 1, 2))(q, k, v)
g2 = jax.grad(loss_ref, argnums=(0, 1, 2))(q, k, v)
for name, a, b in zip("qkv", g1, g2):
check(f"d{name}", a, b, tol=1e-5)import functools
def causal_flash_kernel(q_ref, k_ref, v_ref, o_ref, *, block_q, block_kv):
qi = pl.program_id(0)
q = q_ref[...].astype(jnp.float32)
row0 = qi * block_q
def step(j, state):
m, l, acc = state
kb = jax.lax.dynamic_slice_in_dim(k_ref[...], j * block_kv, block_kv).astype(jnp.float32)
vb = jax.lax.dynamic_slice_in_dim(v_ref[...], j * block_kv, block_kv).astype(jnp.float32)
s = q @ kb.T
cols = j * block_kv + jax.lax.broadcasted_iota(jnp.int32, s.shape, 1)
rows = row0 + jax.lax.broadcasted_iota(jnp.int32, s.shape, 0)
s = jnp.where(cols <= rows, s, -jnp.inf) # edge block: mask inside
m_new = jnp.maximum(m, jnp.max(s, axis=-1))
alpha = jnp.exp(m - m_new)
p = jnp.exp(s - m_new[:, None])
l_new = l * alpha + jnp.sum(p, axis=-1)
acc_new = acc * alpha[:, None] + p @ vb
return m_new, l_new, acc_new
# the causal structure IS the loop bound: blocks past the diagonal never run
n_live = (row0 + block_q + block_kv - 1) // block_kv
m0 = jnp.full((block_q,), -jnp.inf, jnp.float32)
l0 = jnp.zeros((block_q,), jnp.float32)
acc0 = jnp.zeros((block_q, v_ref.shape[1]), jnp.float32)
m, l, acc = jax.lax.fori_loop(0, n_live, step, (m0, l0, acc0))
o_ref[...] = (acc / l[:, None]).astype(o_ref.dtype)def causal_flash(q, k, v, block_q=64, block_kv=64):
sq, d = q.shape
return pl.pallas_call(
functools.partial(causal_flash_kernel, block_q=block_q, block_kv=block_kv),
grid=(sq // block_q,),
in_specs=[
pl.BlockSpec((block_q, d), lambda i: (i, 0)),
pl.BlockSpec(k.shape, lambda i: (0, 0)),
pl.BlockSpec(v.shape, lambda i: (0, 0)),
],
out_specs=pl.BlockSpec((block_q, d), lambda i: (i, 0)),
out_shape=jax.ShapeDtypeStruct(q.shape, q.dtype),
interpret=INTERP,
)(q, k, v)
n = 512
q = jax.random.normal(jax.random.key(0), (n, 64))
k = jax.random.normal(jax.random.key(1), (n, 64))
v = jax.random.normal(jax.random.key(2), (n, 64))
mask = jnp.tril(jnp.ones((n, n), bool))
ref = jax.nn.softmax(jnp.where(mask, q @ k.T, -jnp.inf), axis=-1) @ v
check("causal flash", causal_flash(q, k, v), ref, tol=1e-4)Can you
The gate
| criterion | measured |
|---|---|
| Own flash forward within 1.3x of reference implementations at seq 8192 | est. within 1.3x (pending run) |
| Differential tests green: 1e-3 forward, 1e-2 grads, bf16 | cpu interpret: green at 1e-4 fwd / 1e-5 grads f32; bf16 on-chip pending |
Public artifact for this stage: flash attention derived, and kernels where the data shapes the loop.