Attention, step by step
Code version: DeepInfer @ 9bd9cf2 · 2026-10-21
In one sentence
Attention lets each token take information from earlier tokens. It compares its query with their keys, turns the scores into weights, and mixes their values.
The idea#
The problem: a token alone has no context#
After the embedding, each token has its own vector. The vector of "bank" is the same in "river bank" and in "bank account". The model must mix in information from the other tokens to know which meaning is correct. Attention is the step that does this mix. It is the only place in a transformer where tokens see each other.
Query, key and value#
Each token makes three vectors from its hidden state, with three learned weight matrices (query, key, value):
- The query (q) says what this token looks for.
- The key (k) says what this token offers to others.
- The value (v) is the information that this token gives when another token selects it.
Four steps#
For one token with query q, attention does four steps:
- Score. Compare q with the key of every token: the dot product q·k. A high score means a good match. Divide each score by √dh, where dh is the size of one head's vectors. The deep dive explains why.
- Mask. A token may look only at itself and at earlier tokens. The causal mask sets the scores of later tokens to −∞.
- Softmax. Softmax turns the scores into weights: exp(score) divided by the sum of all exp values. The weights are positive and add up to 1. A masked token gets weight 0, because exp(−∞) = 0.
- Mix. The output is the weighted sum of the values: Σ weightj · vj.
A worked example#
Here is one head with vectors of 2 numbers (dh = 2). The text is "The cat sat on", and we compute the row of "on". These are the same numbers as on the page The KV cache.
| Token | Key | Value | q·k with q = (1, 2) | Score / √2 | Weight |
|---|---|---|---|---|---|
| The | (1, 0) | (2, 0) | 1 | 0.707 | 0.065 |
| cat | (0, 1) | (0, 2) | 2 | 1.414 | 0.131 |
| sat | (1, 1) | (1, 1) | 3 | 2.121 | 0.266 |
| on | (2, 1) | (0, 3) | 4 | 2.828 | 0.539 |
The output of "on" is 0.065·(2, 0) + 0.131·(0, 2) + 0.266·(1, 1) + 0.539·(0, 3) = (0.395, 2.144). The token "on" takes most of its new information from itself and from "sat".
Many heads#
One attention computation is one head. A layer runs many heads in parallel, often 32 to 128. Each head has its own weight matrices, so each head can look for a different relation. For example, one head can find the subject of a verb, and another the previous word. The outputs of all heads are joined and multiplied by one more matrix. This is multi-head attention.
An everyday comparison#
Think of a search in a library. Your question is the query. The label on each book is its key. The text in the book is its value. You compare your question with every label, and you read the books with the best labels.
The comparison stops here. Attention does not select one book. It reads all earlier books and mixes them, with more weight for the better matches. Also, nobody writes the labels: training learns how each token makes its query, key and value.
The picture#
Read panel 2 from left to right. Each row of the grid is one query token, and each column is one key token. The grey cells above the diagonal are later tokens, so the mask hides them. Panel 3 shows the KV heads that the KV cache must store.
Try it: the attention toy#
The toy below has 6 tokens and one head with dh = 2. Tokens 1 to 4 have the keys and values of the worked example, and row "on" has its query q = (1, 2). The queries of the other rows are made up for the toy.
This interactive visual needs JavaScript. Turn on JavaScript in your browser to use it. The text on this page explains the same idea.
Things to try:
- Click each token from left to right. Each row has one more non-zero weight than the row before. This is the causal mask.
- Move Sharpness to 0. Every visible key gets the same weight. Move it to 4. The best key takes almost all the weight.
- Select row "on". Edit the key of "cat" and set k[1] to 3. Watch the weight of "cat" grow and the output move toward the value of "cat", (0, 2).
- Turn off the causal mask. Row "The" can now see "mat". A decoder LLM never does this; an encoder model, such as a model for text classification, can.
In DeepInfer#
Attention is in v4_modules/attn.py. The page MLA attention: one latent vector per token explains the full design. Here are three pointers.
Step 1: the scale is 1/√head_dim#
self.qk_rope_head_dim = cfg.qk_rope_head_dim
self.qk_head_dim = cfg.head_dim
self.q_lora_rank = cfg.q_lora_rank
self.num_heads = cfg.num_attention_heads
self.n_wv_group = cfg.o_groups
self.wv_group_dim = (self.num_heads // self.n_wv_group) * self.qk_head_dim
self.o_head_dim = cfg.o_lora_rank
self.max_model_len = cfg.runtime_max_model_len
self.softmax_scale = self.qk_head_dim**-0.5softmax_scale = qk_head_dim ** -0.5. With head_dim 512 (inferred) the scale is 1/√512 ≈ 0.0442.Step 2: prefill and decode use different kernels#
if input_metadata.is_prompt:
context = self._prefill_attention(
q_b,
kv,
kv_ca_index,
kv_swa_full_index if self.use_swa else None,
total_k_swa_cached,
total_k_swa_cur,
input_metadata,
)
del kv
else:
context = self._decode_attention(
q_b,
kv_ca_index,
kv_swa_index if self.use_swa else None,
kv_cache,
input_metadata,
)Step 3: an index list instead of a mask#
The attention backend calls the FlashMLA library. Its docstring gives the shapes (deepinfer/model_executor/layers/attention_backend.py:35-49):
query:[total_q, num_heads_q, head_dim], BF16. One row per new token, all heads.kv_a_with_mqa:[total_k, 1, head_dim]. The "1" is the single shared KV head.index_topk:[total_q, topk_tokens], int32. For each query, the key rows that it may read.
The plain decode call turns off the causal flag, because the index list already contains only allowed keys. On GPU with 64 or 128 heads, sparse layers and layers without a compressed cache use a fused FlashMLA kernel instead (deepinfer/model_executor/models/v4_modules/attn.py:87-92). That call has no causal flag at all, and it also reads only the listed rows (deepinfer/model_executor/layers/attention_backend.py:107-121).
output, lse = flash_mla_with_kvcache(
q=query,
k_cache=key_cache,
block_table=None, # FlashMLA 保留形参但当前忽略
cache_seqlens=None, # 同上
head_dim_v=head_dim_v,
tile_scheduler_metadata=sched_meta,
softmax_scale=softmax_scale,
causal=False, # sparse attention 只支持非 causal
is_fp8_kvcache=True, # V4.1 的 KV cache 为 paged FP8
indices=indices_in_kvcache.unsqueeze(1),
attn_sink=attn_sink,Translated comments (Chinese → English)
- 128 FlashMLA keeps this parameter but ignores it at present.
- 129 The same.
- 133 Sparse attention supports only non-causal mode.
- 134 The V4.1 KV cache is paged FP8.
Deep dive#
Deep diveThe full formula with shapes
Take one head with T new query tokens and S key tokens. S ≥ T; in a decode step, S is the cached length plus the new tokens.
Q : [T, d_h] K : [S, d_h] V : [S, d_v]
S_raw = Q · Kᵀ [T, S]
S = S_raw · scale scale = 1 / sqrt(d_h)
S = S + M M[i, j] = 0 if key j is allowed for query i, else −∞
P = softmax(S, along each row) [T, S], each row sums to 1
O = P · V [T, d_v]
With H heads, each head does this with its own Q, K and V. The outputs [T, H, d_v] are joined to [T, H · d_v] and multiplied by an output matrix to get back to [T, d].
Cost. The two matrix products cost about 2·T·S·dh + 2·T·S·dv FLOPs per head. They grow with T·S, so a long prompt costs much more than a short one. Double the prompt, and the attention work of a prefill becomes about four times as large.
Memory of the score matrix. Illustrative numbers: T = S = 8,192, H = 128, float32. The full matrix P needs 8,192 × 8,192 × 128 × 4 bytes ≈ 34 GB. No engine stores it. Fused attention kernels compute P in tiles and keep only a running maximum and a running sum per row. FlashMLA returns that per-row value as lse (log-sum-exp), shape [total_q, num_heads_q], fp32 (deepinfer/model_executor/layers/attention_backend.py:46-48).
Deep diveWhy divide by √d_h?
Take q and k with dh random numbers each, every number with mean 0 and variance 1. Then q·k is a sum of dh products. Each product has variance 1, so the sum has variance dh and a typical size of √dh.
- dh = 2: typical score about 1.4.
- dh = 512 (DeepInfer, inferred): typical score about 22.6.
With scores of size 20 or more, softmax puts almost all the weight on one token. The gradients of the other tokens then become almost zero, and training gets stuck. Division by √dh brings the typical score back to about 1.
The effect in the worked example. Row "on", scores 1, 2, 3, 4:
| The | cat | sat | on | |
|---|---|---|---|---|
| softmax(scores), no scale | 0.032 | 0.087 | 0.237 | 0.644 |
| softmax(scores / √2) | 0.065 | 0.131 | 0.266 | 0.539 |
The scaled weights are flatter. In the toy above, Sharpness m multiplies the scaled scores, and m = 1 is the real rule. A value of m = √2 ≈ 1.41 would give the "no scale" row. The slider moves in steps of 0.25, so m = 1.5 is the nearest setting (weights 0.028, 0.080, 0.230, 0.663).
Deep diveSoftmax numerics
Subtract the maximum. exp(z) overflows float32 for z above about 88.7. Kernels therefore compute softmax as:
m = max_j z_j
p_j = exp(z_j − m) / Σ_k exp(z_k − m)
The result is the same, because the factor exp(−m) cancels. Now the largest exponent is exp(0) = 1, so nothing overflows.
Worked example. Scores (1000, 1001). exp(1000) overflows to infinity, and inf / inf is NaN. With the maximum subtracted: exp(−1) = 0.368 and exp(0) = 1, so the weights are 0.269 and 0.731.
Masked entries. −∞ − m is still −∞, and exp(−∞) = 0, so a masked key gets exactly weight 0. A row with only masked entries gives 0 / 0 = NaN. The causal mask never does this, because each token may always see itself.
Log-sum-exp. The value lse = m + ln Σk exp(zk − m) summarizes a row. With it, two partial results of the same row (for example, from two groups of keys) can be joined exactly later. This is why attention kernels often return it.
The toy uses the same rule. The script subtracts the row maximum before exp and gives masked cells weight 0 (assets/viz/attention_toy.js).
Deep diveThe causal mask in prefill and in decode
Prefill. All T prompt tokens are queries at the same time. Query i may read keys 0 … i. The mask is a lower triangle: T·(T + 1)/2 allowed pairs out of T². For T = 4 that is 10 of 16, as in the figure.
Decode. The new token is the last one. It may read every cached key and itself, so a decode step needs no triangle. With draft tokens, a decode step has several new tokens. Each new token may read the cache and the new tokens before it, so the new rows need a small triangle among themselves.
In DeepInfer the "mask" is a list of indices. For each new token, the batch metadata holds the key rows that it may read:
kv_swa_full_index/kv_swa_index: shape[total_q, window_size], the rows of the sliding window (deepinfer/model_executor/models/v4_modules/attn.py:853,deepinfer/model_executor/models/v4_modules/attn.py:857).kv_ca_index: rows of the compressed cache. In sparse layers, an indexer selects the top-k rows (deepinfer/model_executor/models/v4_modules/attn.py:802-809). See Sparse attention with an indexer.
A token never gets an index of a later token, so the index list does the work of the causal mask. That is why the plain decode call passes causal=False (deepinfer/model_executor/layers/attention_backend.py:133), and the fused call has no causal argument (deepinfer/model_executor/layers/attention_backend.py:107-121).
Deep diveHeads, MQA, GQA, MLA and KV memory
The KV cache stores, for each token and layer, one key and one value per KV head:
bytes per token per layer = 2 × n_kv_heads × d_h × bytes_per_number
Illustrative model (not DeepInfer): 32 query heads, dh = 128, BF16 (2 bytes).
| Design | KV heads | Bytes per token per layer | Relative |
|---|---|---|---|
| MHA (multi-head) | 32 | 2 × 32 × 128 × 2 = 16,384 | 1 |
| GQA, groups of 4 | 8 | 2 × 8 × 128 × 2 = 4,096 | 1/4 |
| MQA | 1 | 2 × 1 × 128 × 2 = 512 | 1/32 |
Decode steps read the whole cache for every new token, so a smaller cache also makes decode faster. The cost: with fewer KV heads, the heads can no longer look at different things. This can lower quality.
MLA (multi-head latent attention) keeps one small latent vector per token and lets each query head use it through its own learned projections. In DeepInfer, one cache entry has 512 numbers and serves all heads. Layers with KV compression keep one entry for a group of tokens; the ratio of each layer comes from the checkpoint config (deepinfer/hugging_face/config.py:148-152). With the default FP4 format and its scale bytes, one entry takes 288 bytes. The FP8 option takes 528 bytes (deepinfer/model_executor/layers/tilelang_ops/cache_common.py:109-128, deepinfer/hugging_face/config.py:302-303). The code checks one KV head (deepinfer/worker/cache_engine.py:624). The page MLA attention: one latent vector per token gives the full design; the page The KV cache compares the bytes.
What can go wrong#
- No scale. Without the division by √dh, softmax becomes very sharp for large heads (deep dive "Why divide by √d_h?"). DeepInfer passes the scale to every attention call (
deepinfer/model_executor/models/v4_modules/attn.py:621,deepinfer/model_executor/models/v4_modules/attn.py:664). - A wrong mask. If a query can read a later token, the model sees text that it must predict. In DeepInfer, a wrong index list would do the same, because the kernel does not add its own causal mask.
- Overflow. Softmax without the maximum subtracted gives infinity or NaN for large scores.
- Results that depend on the batch. The same request can give slightly different numbers in different batches, because the kernels sum in a different order. DeepInfer has a batch-invariant mode,
DEEPINFER_BATCH_INVARIANT, that it passes to the decode kernel (deepinfer/model_executor/models/v4_modules/attn.py:666). See Determinism and batch invariance.
Why it matters. Attention is the only part of the model whose cost grows with the length of the text. It is also the reason for the KV cache. Most of DeepInfer's model pages (MLA, sliding window, sparse indexer, KV compression) are ways to make attention cheaper in time or in memory.
Check yourself#
- 1In the worked example, why does the row of "The" have only one non-zero weight?easy
Show answer
"The" is the first token. The causal mask hides all later tokens, so the only allowed key is its own. Softmax over one score gives weight 1.
- 2Compute the weights for the scores (2, 2, 0) after the division by √dh has been applied.medium
Show answer
exp(2) = 7.389, exp(2) = 7.389, exp(0) = 1. Sum: 15.778. Weights: 0.468, 0.468, 0.063.
- 3An illustrative model has 64 query heads, dh = 128 and BF16. How many KV cache bytes per token per layer does it need with MHA, and with GQA in groups of 8?medium
Show answer
MHA: 2 × 64 × 128 × 2 = 32,768 bytes. GQA with groups of 8 has 8 KV heads: 2 × 8 × 128 × 2 = 4,096 bytes, one eighth.
- 4DeepInfer's plain decode attention call passeshard
causal=False. Why does a decode token still not see later tokens?Show answer
The kernel reads only the key rows in the index lists that the batch metadata gives, for example
kv_swa_indexof shape[total_q, window_size](deepinfer/model_executor/models/v4_modules/attn.py:857). These lists never contain a later position. The translated comment says that sparse attention supports only non-causal mode (deepinfer/model_executor/layers/attention_backend.py:133). The fused decode kernel has no causal argument, so it depends on the same lists (deepinfer/model_executor/layers/attention_backend.py:107-121). - 5With dh = 512 and random unit-variance q and k, what is the typical size of q·k, and what does DeepInfer multiply it by?hard
Show answer
The variance is 512, so the typical size is √512 ≈ 22.6. DeepInfer multiplies by
softmax_scale = head_dim ** -0.5= 1/√512 ≈ 0.0442 (deepinfer/model_executor/models/v4_modules/attn.py:78), which brings the typical size back to about 1.