DeepInfer Explained

Attention, step by step

BeginnerLearn the basicsabout 9 min read + 5 min deep dives

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:

  1. 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.
  2. Mask. A token may look only at itself and at earlier tokens. The causal mask sets the scores of later tokens to −∞.
  3. 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.
  4. 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.

TokenKeyValueq·k with q = (1, 2)Score / √2Weight
The(1, 0)(2, 0)10.7070.065
cat(0, 1)(0, 2)21.4140.131
sat(1, 1)(1, 1)32.1210.266
on(2, 1)(0, 3)42.8280.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#

Figure: where Q, K and V come from (panel 1), the four steps for one row (panel 2), and why fewer KV heads save memory (panel 3). Click to zoom.
Open full size ↗

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.

InteractiveAttention toy: scores, mask, softmax, mix

This interactive visual needs JavaScript. Turn on JavaScript in your browser to use it. The text on this page explains the same idea.

Keyboard: Tab to a token or a control; press Enter or Space on a token. Use the arrow keys on a slider. The values are fixed, so every reader sees the same example; Reset starts again.

Things to try:

  1. Click each token from left to right. Each row has one more non-zero weight than the row before. This is the causal mask.
  2. Move Sharpness to 0. Every visible key gets the same weight. Move it to 4. The best key takes almost all the weight.
  3. 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).
  4. 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.5
softmax_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,
            )
A prefill step computes attention for many new tokens at once. A decode step reads the keys and values from the cache. The page Prefill and decode explains the two phases.

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 dive

The 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 dive

Why 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:

Thecatsaton
softmax(scores), no scale0.0320.0870.2370.644
softmax(scores / √2)0.0650.1310.2660.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 dive

Softmax 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 dive

The 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:

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 dive

Heads, 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).

DesignKV headsBytes per token per layerRelative
MHA (multi-head)322 × 32 × 128 × 2 = 16,3841
GQA, groups of 482 × 8 × 128 × 2 = 4,0961/4
MQA12 × 1 × 128 × 2 = 5121/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#

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#

  1. 1
    In 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.

  2. 2
    Compute 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.

  3. 3
    An 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.

  4. 4
    DeepInfer's plain decode attention call passes causal=False. Why does a decode token still not see later tokens?
    hard
    Show answer

    The kernel reads only the key rows in the index lists that the batch metadata gives, for example kv_swa_index of 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).

  5. 5
    With 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.