vllm.v1.attention.backends.mla.sparse_swa ¶

class ComputePrefillMetadataKernel(
    VllmJitKernel["ComputePrefillMetadataKernel.CompileKey"]
):
    @dataclass(frozen=True)
    class CompileKey:
        BLOCK_SIZE: int

    @staticmethod
    @triton.jit(do_not_specialize=["num_prefills", "num_decodes", "window_size"])
    def kernel(
        # Outputs
        prefill_gather_lens_ptr,
        # Inputs
        seq_lens_ptr,
        query_start_loc_ptr,
        num_prefills,
        num_decodes,
        window_size,
        BLOCK_SIZE: tl.constexpr,
    ):
        """Compute prefill gather_lens in a single pass."""
        offset = tl.arange(0, BLOCK_SIZE)
        mask = offset < num_prefills
        # SM12x + Triton 3.6 raises IMA on out-of-bounds address arithmetic for
        # masked-off lanes even though the load mask gates the actual read, so
        # clamp the offset. Caller guarantees num_prefills > 0.
        safe_offset = tl.minimum(offset, num_prefills - 1)

        seq_len = tl.load(seq_lens_ptr + num_decodes + safe_offset, mask=mask)
        qsl_start = tl.load(query_start_loc_ptr + num_decodes + safe_offset, mask=mask)
        qsl_end = tl.load(
            query_start_loc_ptr + num_decodes + safe_offset + 1, mask=mask
        )

        query_len = qsl_end - qsl_start
        prefix_len = seq_len - query_len
        gather_len = query_len + tl.minimum(prefix_len, window_size - 1)

        tl.store(prefill_gather_lens_ptr + offset, gather_len, mask=mask)

    def dispatch(  # type: ignore[override]
        self,
        *,
        num_prefills: int,
    ) -> CompileKey:
        return self.CompileKey(
            BLOCK_SIZE=next_power_of_2(num_prefills),
        )

    def get_warmup_keys(self, vllm_config: VllmConfig) -> list[CompileKey]:
        scheduler_config = vllm_config.scheduler_config
        max_prefills = max(
            1,
            min(
                scheduler_config.max_num_seqs,
                scheduler_config.max_num_batched_tokens,
            ),
        )
        return self._trace_dispatch(self.dispatch)(
            num_prefills=WarmupIntRange(1, max_prefills + 1),
        )

    def compile(self, compile_key: CompileKey) -> None:
        warmup = getattr(self.kernel, "warmup", None)
        assert warmup is not None
        int32_ptr = TritonWarmupTensor(torch.int32)
        warmup(
            int32_ptr,
            int32_ptr,
            int32_ptr,
            compile_key.BLOCK_SIZE,
            0,
            1,
            BLOCK_SIZE=compile_key.BLOCK_SIZE,
            grid=(1,),
        )

    def __call__(
        self,
        prefill_gather_lens: torch.Tensor,
        seq_lens: torch.Tensor,
        query_start_loc: torch.Tensor,
        num_prefills: int,
        num_decodes: int,
        window_size: int,
    ) -> None:
        compile_key = self.dispatch(num_prefills=num_prefills)
        self.kernel[(1,)](
            prefill_gather_lens,
            seq_lens,
            query_start_loc,
            num_prefills,
            num_decodes,
            window_size,
            BLOCK_SIZE=compile_key.BLOCK_SIZE,
        )

class DeepseekSparseSWAMetadataBuilder(AttentionMetadataBuilder):
    """Builds metadata for DeepseekV4 SWA cache.

    Similar to the indexer, this handles mixed batches by:
    1. Using split_decodes_and_prefills() to determine the boundary
    2. Building separate metadata for decode and prefill portions

    Supports:
    - Mixed decode/prefill batches
    - MTP (Multi-Token Prediction) where decode has query_len > 1
    - Chunked prefill (aligns with the indexer's chunking)
    """

    reorder_batch_threshold: int | None = None
    _cudagraph_support: ClassVar[AttentionCGSupport] = AttentionCGSupport.UNIFORM_BATCH
    supports_draft_decode_metadata_update = True

    def __init__(self, *args, **kwargs):
        super().__init__(*args, **kwargs)
        assert isinstance(self.kv_cache_spec, SlidingWindowMLASpec | MLAAttentionSpec)
        mla_spec = cast(SlidingWindowMLASpec | MLAAttentionSpec, self.kv_cache_spec)
        self.head_size = mla_spec.head_size  # Already considered quantization.
        assert isinstance(mla_spec.tokens_per_state, int)
        self.compress_ratio = mla_spec.tokens_per_state
        self.block_size = mla_spec.block_size
        self.max_model_len = self.vllm_config.model_config.max_model_len
        self.max_num_batched_tokens = (
            self.vllm_config.scheduler_config.max_num_batched_tokens
        )

        # Handle MTP: adjust decode_threshold like the indexer does
        spec_config = self.vllm_config.speculative_config
        self.num_speculative_tokens = (
            spec_config.num_speculative_tokens if spec_config else 0
        )
        # Decode can have query_len up to
        #   1 + (2 if parallel drafting else 1) * num_speculative_tokens.
        # sparse_swa has no MQA-vs-dense-MHA routing, so multi-token queries take
        # the prefill path and the decode/prefill split stays at that width.
        spec_mult = (
            2 if (spec_config is not None and spec_config.parallel_drafting) else 1
        )
        self.decode_threshold = 1 + spec_mult * self.num_speculative_tokens
        self.reorder_batch_threshold = None

        hf_config = self.vllm_config.model_config.hf_config
        assert hasattr(hf_config, "sliding_window")
        self.window_size = hf_config.sliding_window

        # Detect which DeepseekV4 layer types this model uses so we only build a
        # FlashMLA tile-scheduler plan for types that will actually be called.
        # Models without compress_ratios (pure SWA) fall back to swaonly.
        compress_ratios = getattr(hf_config, "compress_ratios", None) or [1]
        self._layer_types: set[str] = set()
        for ratio in compress_ratios:
            self._layer_types.add(_layer_type_for(int(ratio)))

        max_tokens = self.vllm_config.scheduler_config.max_num_batched_tokens
        self.token_to_req_indices = torch.zeros(
            max_tokens,
            dtype=torch.int32,
            device=self.device,
        )
        self.decode_swa_indices = torch.zeros(
            max_tokens,
            1,
            self.window_size,
            dtype=torch.int32,
            device=self.device,
        )
        self.decode_swa_lens = torch.zeros(
            max_tokens,
            dtype=torch.int32,
            device=self.device,
        )
        # Allocated unconditionally — consumer picks paged-direct vs dequant
        # at call time.
        self.prefill_swa_indices = torch.zeros(
            max_tokens,
            1,
            self.window_size,
            dtype=torch.int32,
            device=self.device,
        )
        self.prefill_swa_lens = torch.zeros(
            max_tokens,
            dtype=torch.int32,
            device=self.device,
        )
        self.is_valid_token = torch.zeros(
            max_tokens,
            dtype=torch.bool,
            device=self.device,
        )

        # DSpark draft: the block is non-causal (every query attends to the
        # trailing window of context PLUS all query tokens, including future ones),
        # so its per-token index list is wider than `window_size`. The kernel pads
        # the q-head count to B_TOPK. Pad to a kernel-supported width; the logical
        # SWA window remains unchanged when the padded matrix is built.
        self.is_dspark = spec_config is not None and spec_config.use_dspark()
        self.noncausal_index_width = (
            get_dspark_swa_index_width(
                self.window_size,
                self.num_speculative_tokens,
            )
            if self.is_dspark
            else 0
        )
        self.decode_swa_indices_noncausal: torch.Tensor | None = None
        self._max_tokens = max_tokens

    def build(
        self,
        common_prefix_len: int,
        common_attn_metadata: CommonAttentionMetadata,
        fast_build: bool = False,
    ) -> DeepseekSparseSWAMetadata:
        """Build SWA metadata for mixed decode/prefill batches.

        The batch is assumed to be reordered with decodes first (by vLLM scheduler).
        We use split_decodes_and_prefills() to find the boundary, then build
        separate window_topk_idxs for each portion.

        For prefill, we use chunked prefill to align with the indexer's chunking.
        """
        seq_lens = common_attn_metadata.seq_lens
        seq_lens_cpu = common_attn_metadata.seq_lens_cpu_upper_bound
        query_start_loc = common_attn_metadata.query_start_loc
        query_start_loc_cpu = common_attn_metadata.query_start_loc_cpu
        block_table = common_attn_metadata.block_table_tensor
        slot_mapping = common_attn_metadata.slot_mapping

        # Split into decode and prefill portions using configurable threshold
        (num_decodes, num_prefills, num_decode_tokens, num_prefill_tokens) = (
            split_decodes_and_prefills(
                common_attn_metadata, decode_threshold=self.decode_threshold
            )
        )

        # NOTE: Ensure all metadata tensors maintain fixed memory addresses
        # for CUDA graph compatibility.
        token_to_req_indices = common_attn_metadata.token_to_req_indices(
            self.token_to_req_indices
        )

        is_valid_token = self.is_valid_token[: slot_mapping.shape[0]]
        is_valid_token.copy_(slot_mapping >= 0)

        non_causal = not common_attn_metadata.causal
        decode_swa_width = (
            self.noncausal_index_width if non_causal else self.window_size
        )
        decode_swa_indices = self.decode_swa_indices
        if num_decode_tokens > 0:
            self.decode_swa_lens[num_decode_tokens:] = 0
            if non_causal:
                assert self.is_dspark, (
                    "Non-causal DeepseekV4 SWA is only supported for the DSpark "
                    "speculation mode, but causal=False was set without DSpark."
                )
                if self.decode_swa_indices_noncausal is None:
                    self.decode_swa_indices_noncausal = torch.zeros(
                        self._max_tokens,
                        1,
                        self.noncausal_index_width,
                        dtype=torch.int32,
                        device=self.device,
                    )
                decode_swa_indices = self.decode_swa_indices_noncausal
                _compute_dspark_noncausal_swa_indices_kernel[(num_decode_tokens,)](
                    decode_swa_indices,
                    decode_swa_indices.stride(0),
                    self.decode_swa_lens,
                    self.window_size,
                    self.noncausal_index_width,
                    query_start_loc,
                    seq_lens,
                    token_to_req_indices,
                    is_valid_token,
                    block_table,
                    block_table.stride(0),
                    self.block_size,
                    token_offset=0,
                    TRITON_BLOCK_SIZE=1024,
                )
            else:
                _compute_swa_indices_and_lens_kernel[(num_decode_tokens,)](
                    decode_swa_indices,
                    decode_swa_indices.stride(0),
                    self.decode_swa_lens,
                    self.window_size,
                    query_start_loc,
                    seq_lens,
                    token_to_req_indices,
                    is_valid_token,
                    block_table,
                    block_table.stride(0),
                    self.block_size,
                    token_offset=0,
                    TRITON_BLOCK_SIZE=1024,
                )

        # Prefill SWA indices live in paged coordinates. `token_offset` lets
        # the kernel read is_valid_token / token_to_req_indices at absolute
        # prefill positions while writing output starting at index 0.
        if num_prefill_tokens > 0:
            prefill_swa_indices = self.prefill_swa_indices[:num_prefill_tokens]
            prefill_swa_lens = self.prefill_swa_lens[:num_prefill_tokens]
            _compute_swa_indices_and_lens_kernel[(num_prefill_tokens,)](
                prefill_swa_indices,
                prefill_swa_indices.stride(0),
                prefill_swa_lens,
                self.window_size,
                query_start_loc,
                seq_lens,
                token_to_req_indices,
                is_valid_token,
                block_table,
                block_table.stride(0),
                self.block_size,
                token_offset=num_decode_tokens,
                TRITON_BLOCK_SIZE=1024,
            )

        # Pre-compute DeepseekV4 prefill metadata shared across all attention layers.
        deepseek_v4_fields = self._build_deepseek_v4_metadata(
            num_decodes,
            num_prefills,
            seq_lens,
            seq_lens_cpu,
            query_start_loc,
            query_start_loc_cpu,
        )

        # Per-layer-type tile-scheduler plan holders. Empty FlashMLASchedMeta
        # per present DeepseekV4 layer type; the first flash_mla_with_kvcache call of
        # each type triggers the planner and all same-type layers reuse the
        # resulting plan for the rest of the step.
        tile_sched = self.build_tile_scheduler(num_decode_tokens)

        return DeepseekSparseSWAMetadata(
            seq_lens=seq_lens,
            query_start_loc=query_start_loc,
            query_start_loc_cpu=query_start_loc_cpu,
            block_table=block_table,
            slot_mapping=slot_mapping,
            is_valid_token=is_valid_token,
            token_to_req_indices=token_to_req_indices,
            decode_swa_indices=decode_swa_indices[:num_decode_tokens],
            decode_swa_lens=self.decode_swa_lens[:num_decode_tokens],
            decode_swa_width=decode_swa_width,
            prefill_swa_indices=(
                self.prefill_swa_indices[:num_prefill_tokens]
                if num_prefill_tokens > 0
                else None
            ),
            prefill_swa_lens=(
                self.prefill_swa_lens[:num_prefill_tokens]
                if num_prefill_tokens > 0
                else None
            ),
            block_size=self.block_size,
            num_decodes=num_decodes,
            num_prefills=num_prefills,
            num_decode_tokens=num_decode_tokens,
            num_prefill_tokens=num_prefill_tokens,
            # Upper bound on decode-split rows for the kernel's max_q_len
            # hint. common max_query_len bounds every row (scheduled max under
            # adaptive verification), clamped to what the split can admit so a
            # mixed batch's prefill max does not inflate decode scheduling.
            max_decode_query_len=min(
                common_attn_metadata.max_query_len, self.decode_threshold
            ),
            tile_sched_swaonly=tile_sched[_LAYER_TYPE_SWAONLY],
            tile_sched_c4a=tile_sched[_LAYER_TYPE_C4A],
            tile_sched_c128a=tile_sched[_LAYER_TYPE_C128A],
            **deepseek_v4_fields,  # type: ignore[arg-type]
        )

    def update_draft_decode_metadata(
        self,
        metadata: DeepseekSparseSWAMetadata,
    ) -> None:
        if metadata.num_decode_tokens == 0:
            return
        assert metadata.query_start_loc is not None
        assert metadata.seq_lens is not None
        assert metadata.token_to_req_indices is not None
        assert metadata.is_valid_token is not None
        assert metadata.decode_swa_indices is not None
        assert metadata.decode_swa_lens is not None

        _compute_swa_indices_and_lens_kernel[(metadata.num_decode_tokens,)](
            metadata.decode_swa_indices,
            metadata.decode_swa_indices.stride(0),
            metadata.decode_swa_lens,
            metadata.decode_swa_indices.shape[-1],
            metadata.query_start_loc,
            metadata.seq_lens,
            metadata.token_to_req_indices,
            metadata.is_valid_token,
            metadata.block_table,
            metadata.block_table.stride(0),
            self.block_size,
            token_offset=0,
            TRITON_BLOCK_SIZE=1024,
        )
        tile_sched = self.build_tile_scheduler(metadata.num_decode_tokens)
        metadata.tile_sched_swaonly = tile_sched[_LAYER_TYPE_SWAONLY]
        metadata.tile_sched_c4a = tile_sched[_LAYER_TYPE_C4A]
        metadata.tile_sched_c128a = tile_sched[_LAYER_TYPE_C128A]
        metadata.flashinfer_sparse_index_cache.clear()

    def build_tile_scheduler(
        self, num_decode_tokens: int
    ) -> dict[str, FlashMLASchedMeta | None]:
        """Allocate one empty ``FlashMLASchedMeta`` per present DeepseekV4 layer type.

        Returned instances have ``tile_scheduler_metadata`` / ``num_splits``
        set to ``None``; the FlashMLA C++ decode path will allocate them and
        run the tile-scheduler planner on the first ``flash_mla_with_kvcache``
        call of each type. Subsequent same-type calls reuse the plan because
        the tensors (and ``have_initialized``) are populated on the struct.

        Returns all-``None`` when there are no decode tokens this step, so
        ``_forward_decode`` sees a clean sentinel.
        """
        out: dict[str, FlashMLASchedMeta | None] = {
            _LAYER_TYPE_SWAONLY: None,
            _LAYER_TYPE_C4A: None,
            _LAYER_TYPE_C128A: None,
        }
        if (
            num_decode_tokens == 0
            or current_platform.is_rocm()
            or current_platform.is_xpu()
            or current_platform.is_device_capability_family(120)
        ):
            return out
        for layer_type in self._layer_types:
            # get_mla_metadata() is the official FlashMLA entry point that
            # returns a fresh empty FlashMLASchedMeta; using it keeps this
            # call site aligned with the rest of the vLLM FlashMLA backends
            # that already go through the same stub.
            out[layer_type] = get_mla_metadata()[0]
        return out

    def _build_deepseek_v4_metadata(
        self,
        num_decodes: int,
        num_prefills: int,
        seq_lens: torch.Tensor,
        seq_lens_cpu: torch.Tensor | None,
        query_start_loc: torch.Tensor,
        query_start_loc_cpu: torch.Tensor,
    ) -> dict[str, torch.Tensor | int | None]:
        """Pre-compute DeepseekV4 prefill metadata during the metadata build phase.

        Returns a dict of keyword arguments to pass to the
        DeepseekSparseSWAMetadata constructor.

        Note: C128A sparse metadata is computed by the FlashMLASparse builder
        (which owns the C128A block_table), not here.
        """
        result: dict[str, torch.Tensor | int | None] = {}

        # --- Prefill query metadata (single Triton kernel + CPU slicing) ---
        if num_prefills > 0:
            assert seq_lens_cpu is not None
            pfx_gather_lens = torch.empty(
                num_prefills, dtype=torch.int32, device=seq_lens.device
            )
            _COMPUTE_PREFILL_METADATA_KERNEL(
                pfx_gather_lens,
                seq_lens,
                query_start_loc,
                num_prefills,
                num_decodes,
                self.window_size,
            )

            result["prefill_seq_lens"] = seq_lens[num_decodes:]
            result["prefill_seq_lens_cpu"] = seq_lens_cpu[num_decodes:]
            result["prefill_gather_lens"] = pfx_gather_lens
            result["prefill_query_lens_cpu"] = (
                query_start_loc_cpu[num_decodes + 1 : num_decodes + num_prefills + 1]
                - query_start_loc_cpu[num_decodes : num_decodes + num_prefills]
            ).to(dtype=torch.int32)
            result["prefill_window_size"] = self.window_size
            result["prefill_max_model_len"] = self.max_model_len
            result["prefill_max_num_batched_tokens"] = self.max_num_batched_tokens

        return result