Sliding-window attention

A decode step writes all its new SWA vectors first, then runs attention. With the default 5 draft tokens, one step has 6 query tokens per sequence: positions p to p + 5 (see Speculative decoding). The rule is in csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:103-104 and deepinfer/model_executor/models/v4_modules/attn.py:421-432.

1. The first query (position p) needs positions p − 127 to p.
2. The step also writes positions p + 1 to p + 5.
3. So the ring must hold positions p − 127 to p + 5 at the same time: 133 positions.
4. With 128 slots, the draft at p + 1 would overwrite position p − 127, which the first query still reads.

The code reserves DEEPINFER_USE_NEXTN_DECODE + 1 = 6 extra slots, so the ring has 134 slots: one more than this count needs. The docs give two reasons for the extra slots. The hit policy calls them "6 tokens of speculative backtracking" (doc/inference_deployment_qa.md:191). The store copy "can overshoot the boundary by at most 5 tokens, which is absorbed by the 6 extra positions" (doc/inference_deployment_qa.md:165).

Rejected drafts. A rejected draft is never deleted. The next step starts at the first rejected position and writes the same positions again, into the same slots. Then the ring holds the correct tokens again.

The compression state uses the same rule with its own numbers: its ring must hold ratio + tokens per step − 1 entries (deepinfer/model_executor/layers/tilelang_ops/compress.py:98-100). See KV compression (C1 and C2).

Bytes. One ring per layer and sequence: 134 × 528 B = 70,752 B. The code pads each layer to 32 bytes; 70,752 is already a multiple of 32 (deepinfer/worker/cache_engine.py:274-277). The SWA format is always MXFP8, also when the main cache uses FP4: use_fp4=False for SWA (deepinfer/worker/cache_engine.py:261-263).

Tensor view. swa_cache is uint8 [num_state_slots, 134, 1, 528]. The row stride is the size of the whole state slot, so one row index (state_id) selects one sequence (deepinfer/worker/cache_engine.py:419-422, deepinfer/worker/cache_engine.py:216-218).

Where in the slot. The slot starts with the "stored" SWA rings of the producer layers. Then come the "dropped" SWA rings of the YOCO consumer layers, and the compression state (deepinfer/worker/cache_engine.py:221-241). Only the stored part goes to the external KV store. The deployment notes give its size: "20 layers × 134 tokens × 528 B = 1,415,040 B per sequence" (doc/inference_deployment_qa.md:169). The number 20 comes from the docs, not from the code. See YOCO and SWA bounded replay.

Sequence with state_id 3, decode query at position 1,000, window 128, ring 134 (csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:644-649):

1. real_window_size = min(1,000 + 1, 128) = 128.
2. The oldest key is 1,000 − 128 + 1 = 873. 873 % 134 = 69, so offset = 3 × 134 + 69 = 471.
3. The kernel makes 128 slots: 402 + (69 + j) % 134 for j = 0 … 127 (deepinfer/model_executor/layers/tilelang_ops/get_kv_index.py:91).
4. j = 0 … 64 give slots 471 … 535 (positions 873 … 937). j = 65 … 127 wrap and give slots 402 … 464 (positions 938 … 1,000).
5. Slot 464 = 3 × 134 + 1,000 % 134 = 402 + 62: the query's own key, written in the same step (csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:104).

A short sequence at position 50 gets real_window_size = 51, so columns 51 … 127 are −1.

Shapes: kv_swa_index_offset_len int32 [T, 2]; kv_swa_index int32 [T, 128] (deepinfer/model_executor/input_metadata/input_metadata.py:313-317). With zero-overhead scheduling, the device kernel prepare_v4_decode_metadata builds the offsets (deepinfer/model_executor/input_metadata/prepare_tensor_input_zos.py:326-367).

The load plan. C++ copies each sequence's history window from its ring into the first section of the buffer:

    for (auto p_seq : args.seq_metadata_list) {
        TORCH_CHECK(p_seq != nullptr);
        TORCH_CHECK_GE(p_seq->seqlen_q, 0);
        TORCH_CHECK_GE(p_seq->seqlen_k, p_seq->seqlen_q);
        TORCH_CHECK_GE(p_seq->swa_valid_start, 0);

        const int32_t history_end = p_seq->seqlen_k - p_seq->seqlen_q;
        const int32_t history_begin = history_end - window_size;
        const int32_t padding_len = std::clamp(p_seq->swa_valid_start - history_begin, 0, window_size);
        const int32_t valid_begin = history_begin + padding_len;
        builder.append_zero_fill(padding_len);
        builder.append_ring_range(p_seq->state_id, valid_begin, history_end);
    }

    auto result = builder.build();
    TORCH_CHECK_EQ(result.total_k, static_cast<int64_t>(args.seq_metadata_list.size()) * window_size);
    return result;

1. Lines 378-379. The history window is positions q_start − W to q_start − 1: the W tokens before the chunk.
2. Lines 380-382. Positions before 0 (or before swa_valid_start) become zero rows: a segment with src_block_id = −1.
3. Line 383. The rest is a ring range. The builder splits it where the ring wraps, so one sequence has at most three segments (csrc/engine/input_metadata/block_load_plan.cpp:37-46, deepinfer/model_executor/input_metadata/input_metadata.py:268-271).
4. Line 387. The section always has num_seqs × W rows.

Example: 2,048 cached tokens, so the history is 1,920 … 2,047. 1,920 % 134 = 44. The plan has two segments: ring rows 44 … 133 (90 tokens), then rows 0 … 37 (38 tokens). A new prompt (q_start = 0) gets one zero segment of 128 rows.

The pivots.

    for (int i_seq = 0; i_seq < n_seq; ++i_seq) {
        auto p_seq = seq_metadata_list[i_seq];

        int32_t sec1_start = i_seq * window_size;
        int32_t sec2_start = n_seq * window_size + i_q;

        int32_t q_end = p_seq->seqlen_k;
        int32_t q_len = p_seq->seqlen_q;
        int32_t q_start = q_end - q_len;
        // zero-state 重算：[0, valid_start) 的 swa state 被清零，可见窗口左边界 clamp 到 valid_start，
        // 避免 attend 到清零 slot。valid_start=0（正常路径）时与原式 start_in_win = min(W, max(j+1, W - q_start)) 完全等价。
        int32_t valid_start = p_seq->swa_valid_start;
        // 下面的代码假设了 valid_start <= q_start。当前能保证可见范围一定比输入范围大。
        TORCH_CHECK_LE(valid_start, q_start);
        for (int32_t j = 0; j < q_len; ++j, ++i_q) {
            int32_t pos = q_start + j; // 该 query token 的绝对位置
            // 可见 key 的绝对位置下界：滑窗下界与 state 有效下界取 max，再 clamp 到 >= 0
            int32_t abs_lo = std::max(pos - window_size + 1, valid_start);
            if (abs_lo < 0)
                abs_lo = 0;

            // section1（历史，从 state buffer 读）：可见区间 [min(abs_lo, q_start), q_start)
            int32_t sec1_lo_abs = std::min(abs_lo, q_start);
            int32_t start_slot = sec1_lo_abs - (q_start - window_size);
            TORCH_CHECK_GE(start_slot, 0) << "swa start_slot out of range";
            TORCH_CHECK_LE(start_slot, window_size) << "swa start_slot out of range";
            p_kv_swa_full_index_pivots[i_q * 4 + 0] = sec1_start + start_slot;
            p_kv_swa_full_index_pivots[i_q * 4 + 1] = window_size - start_slot;

            // section2（本 chunk 新增）：可见区间 [max(abs_lo, q_start), pos]
            int32_t k_lo = std::max(0, abs_lo - q_start);
            TORCH_CHECK_LE(k_lo, j) << "swa section2 empty: query token sees nothing (valid_start too large?)";
            p_kv_swa_full_index_pivots[i_q * 4 + 2] = sec2_start + k_lo;
            p_kv_swa_full_index_pivots[i_q * 4 + 3] = j - k_lo + 1;
        }

1. Line 507. The oldest visible key is pos − W + 1, not below valid_start and 0.
2. Lines 512-517. Section 1: the history part. start_slot is the first visible row of this sequence's 128 history rows; the length is W − start_slot.
3. Lines 520-523. Section 2: the new part, from k_lo to the query itself.

Example, same sequence (2,048 cached, 300 new). Query j = 0 (position 2,048): abs_lo = 1,921, start_slot = 1, so 127 history keys and 1 new key. Query j = 200 (position 2,248): abs_lo = 2,121, start_slot = 128 (empty history part), k_lo = 73, so 128 new keys.

The store. After the load, store_kv_cache writes the last 134 new vectors into the ring (deepinfer/model_executor/models/v4_modules/attn.py:855). The prefill vectors went through an FP8 round trip (cast_back=True). So attention in the prefill sees the same numbers that decode later reads from the ring (deepinfer/model_executor/models/v4_modules/attn.py:412-419).

• RoPE. A compressed layer builds its rotary table with compress_rope_theta; an SWA-only layer uses rope_theta (deepinfer/model_executor/models/v4_modules/mla_attn.py:275-280). The SWA keys of a layer use that layer's table.
• Projection. The SWA vector comes from the same merged GEMM as the query latent (q_a_kv_swa_proj, deepinfer/model_executor/models/v4_modules/attn.py:280-288).
• Sharing. Cross-layer sharing never shares the ring. Each layer stores its own SWA KV (deepinfer/worker/cache_engine.py:406-422). See Cross-layer KV sharing.
• Draft model. The DFlash layers also have sink_mode="swa" and their own rings (deepinfer/hugging_face/config.py:136-143).

Each sequence has a field swa_valid_start (csrc/engine/sequence.hpp:205). The C++ code treats ring positions before it as missing: the load plan zero-fills them and the pivots do not point at them (csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:380-382, csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:501-509).

In this release, no code sets the field above 0. reset_for_recompute sets it to 0 (csrc/engine/sequence.hpp:223-227), and pybind exposes it to Python as a read-write property (csrc/engine/pybind_sequence.cpp:132), but no Python code writes it. The field is a hook for the "encoder SWA bounded replay" of the docs. After a prefix-cache hit without SWA state, that design recomputes only the last 128 tokens (doc/inference_deployment_qa.md:183-185). The KV cache system that would use it is not in this release (README.md:40).