LoRA serving, embeddings, rerankers, pooling

You have fifty customer-specific fine-tunes of Llama-3-8B. Each is a rank-16 LoRA adapter: 41.9 M trainable parameters against the base model's 8.03 B. Fifty of them together are 4.2 GB in bf16 — a quarter of the 16.06 GB base checkpoint. Serving fifty separate vLLM processes would need fifty copies of those 16.06 GB and fifty H100s. Obviously you serve them from one process.

So you start the server with a generous slot count and send traffic. Two things go wrong.

The first is loud:

    def _apply_adapters(self, lora_requests: set[LoRARequest]) -> None:
        loras_map = {
            lora_request.lora_int_id: lora_request
            for lora_request in lora_requests
            if lora_request
        }
        if len(loras_map) > self._adapter_manager.lora_slots:
            raise RuntimeError(
                f"Number of requested LoRAs ({len(loras_map)}) is greater "
                "than the number of GPU LoRA slots "
                f"({self._adapter_manager.lora_slots})."
            )
        for lora in loras_map.values():
            self.add_adapter(lora)

Different adapters partition prefix-cache identity, so fifty adapters can require fifty copies of a common prompt's KV. That does not make every request cold: each adapter can reuse its own retained prefix. Miss rate depends on working-set size, reuse distance, routing and eviction. At fifty 2048-token copies the illustrated footprint is 13.4 GB, not a proof of zero hits.

This chapter is about both halves of multi-tenant LoRA serving — the kernel that makes mixed-adapter batching possible, and the cache partitioning that makes it expensive — and then about the opposite problem: pooling models, where there is no decode loop, no KV growth, and the hard parts of this book simply do not apply.

With stored $W\in\mathbb{R}^{d_{\mathrm{out}}\times d_{\mathrm{in}}}$, $A\in\mathbb{R}^{r\times d_{\mathrm{in}}}$ and $B\in\mathbb{R}^{d_{\mathrm{out}}\times r}$, $W'=W+(\alpha/r)BA$. Row-vector inputs use $xW^\top+(\alpha/r)(xA^\top)B^\top$. Merged adapters can share hardware through ordinary scheduling but require separate merged weights. Unmerged serving shares the base matrix and applies each row's low-rank update.

The naive way to batch that is to group the batch by adapter and run one GEMM per group. With eight adapters in a batch of 64 you get eight GEMMs of 8 rows each — every one of them latency-bound, and the whole point of continuous batching (§1.3) evaporates.

The trick every production engine uses instead: run the base GEMM exactly once for the whole batch, then apply a per-row indexed low-rank correction. The correction is cheap precisely because it never touches a $d \times d$ matrix — it goes through the rank-$r$ bottleneck, so it is [tokens, r]-shaped in the middle. Both engines call this SGMV (Segmented Gather Matrix-Vector multiplication), after the Punica paper.

Read the FLOP counts on that diagram. The rank-16 correction costs 0.78% of the base GEMM for a square projection. That is the whole reason multi-tenant LoRA serving is viable at all, and the next section makes the number general.

Symbols: $d$ hidden size, $L$ layers, $r$ adapter rank, $B$ batch size in tokens, $k$ the number of distinct adapters active in one batch, $P(r)$ the parameter count of one adapter. Llama-3-8B: $d = 4096$, $L = 32$, 8 KV heads of 128 (so $d_\text{kv} = 1024$), FFN intermediate 14336, 8.03 B parameters, 16.06 GB in bf16.

Adapter size

An adapter on one linear layer costs $r(d_\text{in} + d_\text{out})$ parameters. Summing the seven usual targets for one Llama-3-8B decoder layer:

In bf16 that is exactly $r \cdot 5$ MiB. The arithmetic is unusually clean:

Attention-only adapters (q,k,v,o) are a third of that: $\Sigma_\text{layer} = 26{,}624$, so $P(16) = 13.63$ M, 26 MiB. This is why "dozens of fine-tunes on one GPU" is a real deployment and not a slide.

What the correction costs, in both regimes

Adapter work and base work need not have the same arithmetic intensity: a base matrix is reused across the whole batch, while an adapter may serve only a few rows. Rank, segmentation, padding and intermediate traffic matter. The following are FLOP and ideal weight-byte ratios, not latency predictions. Using the stated prefill parameter convention gives:

Decode divides by the bytes a step actually streams, which §0.4 puts at 15.01 GB rather than 16.06 GB — the input embedding table is gathered one row at a time, not streamed, so it does not appear in the decode floor. The adapter, by contrast, is streamed in full at $r \cdot 5$ MiB:

at $r = 16$ that is $0.52\%$ of a prefill token and $0.56\%$ of a decode step. But the two regimes multiply those differently, and this asymmetry is the operational core of LoRA serving:

In this one-stream-per-active-adapter byte model, k*r about 286 adds 10% of the base weight traffic. It is not a free-operation threshold: launches, per-adapter reuse, caching, padded ranks and achieved bandwidth can dominate well below it. The table does not prove that a particular latency tax is unavoidable or that kernel engineering cannot improve it.

vLLM: preallocated stacked buffers plus a Punica-style Triton kernel

vLLM replaces every adaptable nn.Module with a wrapper from vllm/lora/layers/, and each wrapper preallocates the whole slot array up front — at max_lora_rank, regardless of what the adapters actually are:

        self.lora_a_stacked = tuple(
            torch.zeros(
                max_loras,
                1,
                lora_a_out_size,
                self.input_size,
                dtype=lora_config.lora_dtype,
                device=self.device,
            )
            for _ in range(self.n_slices)
        )
        self.lora_b_stacked = tuple(
            torch.zeros(
                max_loras,
                1,
                lora_b_out_size,
                lora_config.max_lora_rank,
                dtype=lora_config.lora_dtype,
                device=self.device,
            )
            for _ in range(self.n_slices)
        )

The forward path is exactly Figure 1 — base GEMM, then correction, in place:

    def _apply_sync(
        self, x: torch.Tensor, bias: torch.Tensor | None = None
    ) -> torch.Tensor:
        output = self._get_quant_method().apply(self.base_layer, x, bias)
        return self._apply_lora_to_output(x, output)

_apply_lora_to_output calls into the PunicaWrapperGPU — selected by platform, not by a flag: get_punica_wrapper is a Platform classmethod and the CUDA implementation returns "vllm.lora.punica_wrapper.punica_gpu.PunicaWrapperGPU" unconditionally (vllm/platforms/cuda.py:L560-L562). There is no Triton-vs-CUDA choice to make on the vLLM side; the Triton path is the path. It allocates the rank-sized intermediate and fires the two kernels:

        r = lora_b_stacked[0].size(-1)
        # We set the buffer to be float32 by default, refer to:
        # https://github.com/triton-lang/triton/issues/1387
        # Note: buffer is zeroed inside the shrink op
        buffer = torch.empty(
            (len(output_slices), x.size(0), r), dtype=torch.float32, device=x.device
        )
        add_inputs = kwargs.pop("add_inputs", True)
        self.add_shrink(
            buffer,  # type: ignore
            x,
            lora_a_stacked,
            scale,
            **kwargs,
        )
        self.add_expand(
            y,
            buffer,  # type: ignore
            lora_b_stacked,
            output_slices,
            add_inputs=add_inputs,
            **kwargs,
        )

The shrink launcher documents every shape in its own signature — this is the clearest single statement of the segmented-GEMM contract in either repo:

    inputs: torch.Tensor,  #  shape [num_tokens, hidden_size]
    lora_a_weights: list[torch.Tensor],  # shape [num_loras, lora_rank, hidden_size]
    output_tensor: torch.Tensor,  # shape [num_slices, num_tokens, lora_rank]
    token_lora_mapping: torch.Tensor,  # shape [num_tokens]
    token_indices_sorted_by_lora_ids: torch.Tensor,  # shape [num_tokens]
    num_tokens_per_lora: torch.Tensor,  # shape [max-loras + 1]
    lora_token_start_loc: torch.Tensor,  # shape [max-loras + 2]
    lora_ids: torch.Tensor,  # shape [max-loras + 1]
    no_lora_flag_cpu: torch.Tensor,  # shape [1]
    num_active_loras: torch.Tensor,  # CPU tensor [1], number of active LoRAs

The metadata is a sort and a unique, rebuilt every forward pass:

        # token_indices_sorted_by_lora_ids
        _, token_indices_sorted_by_lora_ids = torch.sort(
            token_lora_mapping, stable=True
        )
        # start gpu transfer
        self.token_indices_sorted_by_lora_ids[:num_tokens].copy_(
            token_indices_sorted_by_lora_ids, non_blocking=True
        )

        # active_lora_ids, num_tokens_per_lora
        lora_ids, num_tokens_per_lora = torch.unique(
            token_lora_mapping, sorted=True, return_counts=True
        )

and the kernel's third grid axis is the adapter, so each CTA knows which slot it serves and which rows of the input it owns:

    slice_id = tl.program_id(axis=1)
    lora_idx = tl.program_id(axis=2)

    lora_id = tl.load(lora_ids + lora_idx)
    if lora_id == -1:
        # Early exit for the no-lora case.
        return

    lora_m_size = tl.load(num_tokens_per_lora + lora_idx)

    cta_m_offset = pid_m * BLOCK_M
    if cta_m_offset >= lora_m_size:
        # Early exit CTA.
        return

    # num rows this CTA should process.
    cta_m_len = min(BLOCK_M, lora_m_size - cta_m_offset)

    # Identify all rows that this CTA should process.
    lora_m_indices_start = tl.load(lora_token_start_loc + lora_idx)
    cta_lora_seq_indices = (
        token_indices_sorted_by_lora_ids + lora_m_indices_start + cta_m_offset
    )
    # Load all relevant row indices.
    offset_m = tl.arange(0, BLOCK_M) % cta_m_len
    ram = tl.load(cta_lora_seq_indices + offset_m)

Defaults matter here and they are conservative:

    max_lora_rank: MaxLoRARanks = 16
    """Max LoRA rank."""
    max_loras: int = Field(default=1, ge=1)
    """Max number of LoRAs in a single batch."""
    fully_sharded_loras: bool = False
    """By default, only half of the LoRA computation is sharded with tensor
    parallelism. Enabling this will use the fully sharded layers. At high
    sequence length, max rank or tensor parallel size, this is likely faster.
    """
    max_cpu_loras: int | None = None
    """Maximum number of LoRAs to store in CPU memory. Must be >= than
    `max_loras`."""

MaxLoRARanks is Literal[1, 8, 16, 32, 64, 128, 256, 320, 512] — the rank is a config-hash factor, so it is quantised to keep the compiled-graph cache small. max_loras=1 as a default means multi-adapter batching is opt-in; you must set --max-loras yourself. And the whole subsystem is behind enable_lora: bool = False (vllm/engine/arg_utils.py:L611-L616) — no --enable-lora, no wrapper substitution, no kernels.

SGLang: chunked SGMV, and a rank-aware kernel

SGLang's default backend names its lineage in the docstring:

class ChunkedSgmvLoRABackend(BaseLoRABackend):
    """
    Chunked LoRA backend using segmented matrix-vector multiplication.

    This backend is largely based on the SGMV (Segmented Gather Matrix-Vector multiplication) algorithm
    introduced in the Punica paper (https://arxiv.org/pdf/2310.18547). One main variation made here is to
    segment the input sequences into fixed-size chunks, which reduces excessive kernel launches especially
    when the LoRA distribution is skewed.
    """

    name = "csgmv"
    supports_prefill_cuda_graph = True

The chunking is the real difference from vLLM's approach. vLLM sorts tokens by adapter and gives each adapter a variable-length row range; SGLang cuts the sorted stream into fixed --max-lora-chunk-size segments (default 16, choices 16/32/64/128), so segment count is predictable and CUDA-graph capture works for prefill too. When one adapter owns 90% of a batch, vLLM's per-adapter CTA count is wildly imbalanced; SGLang's is uniform.

The second difference is more consequential, and it is one line of Triton:

    # Current block computes sequence with batch_id,
    # which starts from row seg_start of x with length seg_len
    w_index = tl.load(weight_indices + pid_s)
    rank = tl.load(lora_ranks + w_index)

    # If rank is 0, this kernel becomes a no-op as the output is always trivially correct.
    if rank == 0:
        return

    # Adjust N dim according to the specific LoRA adapter
    cur_n = tl.minimum(N, rank * NUM_SLICES)

SGLang carries a per-adapter rank into the kernel and masks loads past it. vLLM does not: its launcher takes N, K = lora_a_weights[0].shape[-2:] # K=hidden_size,N=rank (lora_shrink_op.py:L197), and that shape is the buffer shape, allocated at max_lora_rank (lora_a_out_size = lora_config.max_lora_rank, vllm/lora/layers/base_linear.py:L111-L120). Consequence: on vLLM, an $r=8$ adapter served under --max-lora-rank 64 moves the bytes of an $r=64$ adapter. Be precise about what SGLang saves, because it is not everything. Both engines allocate at the fleet maximum — SGLang's --max-lora-rank defaults to None and is inferred from the adapters in --lora-paths (python/sglang/srt/server_args.py:L2925-L2929) — and both launch a grid whose N dimension is cdiv(padded_N, BLOCK_N). What cur_n masks is the tl.load, so SGLang saves the memory transactions past the true rank, which is exactly the term that matters in the bandwidth-bound decode regime, while still paying launch overhead for blocks entirely past cur_n (exercise 4). Going back to the $k \cdot r$ rule: vLLM's effective $r$ is the fleet maximum; SGLang's is the per-adapter actual, on the bandwidth term only.

SGLang's flag surface is broader and its defaults are more aggressive:

    max_loras_per_batch: A[
        int,
        "Maximum number of adapters for a running batch, include base-only request.",
        NS("lora"),
    ] = 8
    lora_eviction_policy: A[
        str,
        Arg(
            help="LoRA adapter eviction policy when memory pool is full. 'lru': Least Recently Used (default, better cache efficiency). 'fifo': First-In-First-Out.",
            choices=["lru", "fifo"],
        ),
        NS("lora"),
    ] = "lru"

Which layers, and can the LM head be adapted?

Both, and yes. vLLM's registry enumerates the wrapper classes it will substitute:

_all_lora_classes: tuple[type[BaseLayerWithLoRA], ...] = (
    VocabParallelEmbeddingWithLoRA,
    ColumnParallelLinearWithLoRA,
    MergedColumnParallelLinearWithLoRA,
    QKVParallelLinearWithLoRA,
    MergedQKVParallelLinearWithLoRA,
    RowParallelLinearWithLoRA,
    ReplicatedLinearWithLoRA,
    LogitsProcessorWithLoRA,
    ColumnParallelLinearWithShardedLoRA,
    QKVParallelLinearWithShardedLoRA,
    MergedColumnParallelLinearWithShardedLoRA,
    MergedColumnParallelLinearVariableSliceWithLoRA,
    MergedQKVParallelLinearWithShardedLoRA,
    RowParallelLinearWithShardedLoRA,
    FusedMoEWithLoRA,
    FusedMoE3DWithLoRA,
)

LogitsProcessorWithLoRA is the LM head; VocabParallelEmbeddingWithLoRA is the input embedding; the two FusedMoE* entries adapt expert weights, which matters for the MoE models of §7.1. SGLang states the same set as a flat list of module suffixes accepted by --lora-target-modules:

SUPPORTED_LORA_TARGET_MODULES = [
    "q_proj",
    "k_proj",
    "v_proj",
    "o_proj",
    "q_a_proj",
    "kv_a_proj_with_mqa",
    "q_b_proj",
    "kv_b_proj",
    "wq_b",
    "wk",
    "weights_proj",
    "gate_proj",
    "up_proj",
    "down_proj",
    "qkv_proj",
    "gate_up_proj",
    "embed_tokens",
    "lm_head",
    # Inkling attention projections (merged q/k/v/r and its row-parallel output).
    "qkvr",
    "wo_ud",
]

Note q_a_proj, kv_a_proj_with_mqa, q_b_proj, kv_b_proj — MLA's factorised projections (§7.2) are adaptable, which is not obvious given how the absorbed-weight trick rearranges them.

The lifecycle: registry, slots, eviction

vLLM's model manager keeps a slot array lora_index_to_id and copies weights into the preallocated buffers on activation:

    def activate_adapter(
        self,
        lora_id: int,
    ) -> bool:
        """Move LoRA into a GPU buffer to be used in the forward pass."""
        if lora_id in self._active_adapters:
            return False
        first_free_slot = next(
            (
                (i, lora_id)
                for i, lora_id in enumerate(self.lora_index_to_id)
                if lora_id is None
            ),
            None,
        )
        if first_free_slot is None:
            raise ValueError("No free lora slots")
        index, _ = first_free_slot
        self._active_adapters[lora_id] = None
        lora_model = self._registered_adapters[lora_id]
        logger.debug(
            "Activating LoRA. int id: %d, slot index: %d", lora_model.id, index
        )

with eviction layered on in the LRU subclass:

    def activate_adapter(
        self,
        lora_id: int,
    ) -> bool:
        if (
            lora_id not in self._active_adapters
            and len(self._active_adapters) >= self.lora_slots
        ):
            self._active_adapters.remove_oldest()
        result = super().activate_adapter(lora_id)
        # We always touch to update the LRU cache order
        self._active_adapters.touch(lora_id)
        return result

SGLang's memory pool does the same job with pinning and a pluggable policy:

        def get_available_buffer_slot():
            # 1. Prioritize empty slots
            for buffer_id in range(self.max_loras_per_batch):
                if self.buffer_id_to_uid[buffer_id] == EMPTY_SLOT:
                    return buffer_id

            # 2. Memory pool is full, need to evict using policy
            candidates = set()

            for buffer_id in range(self.max_loras_per_batch):
                uid = self.buffer_id_to_uid[buffer_id]

                # Skip if this adapter is needed by current batch
                if uid in cur_uids:
                    continue

                # Skip if this adapter is pinned
                if uid is not None:
                    lora_ref = lora_refs.get(uid)
                    if lora_ref and lora_ref.pinned:
                        continue

                candidates.add(uid)

            if not candidates:
                raise ValueError(
                    "No available buffer slots found. Please ensure the number of active (pinned) loras is less than max_loras_per_batch."
                )

Admission is checked before a request is allowed to join a running batch:

    def validate_lora_batch(self, lora_ids: set[Optional[str]]) -> bool:
        """
        Validate if the LoRA IDs in the batch can be loaded into the current LoRA memory pool.
        """
        if len(lora_ids) > self.max_loras_per_batch:
            return False

What loading mid-serving costs. Derived: an $r=16$ adapter is 83.9 MB; at PCIe Gen5 x16's 64 GB/s theoretical unidirectional bandwidth, the host-to-device copy alone is 1.3 ms — on the critical path, blocking a forward pass. If the adapter has to come off disk or from the Hugging Face Hub first, it is seconds. SGLang's --enable-lora-overlap-loading exists exactly to hide that transfer behind compute, at the cost of pinning adapter weights in CPU memory: enabling it makes --max-loaded-loras mandatory and asserts it is at most twice --max-loras-per-batch (python/sglang/srt/server_args.py:L9446-L9456). It is off by default — the flag resolves to False when unset.

SGLang also has something vLLM does not: a drainer, which handles the tail-latency pathology where a few popular adapters monopolise all eight slots and an unpopular tenant starves forever.

    def _can_schedule_lora_req(
        self, req: Req, running_loras: set[Optional[str]]
    ) -> bool:
        """
        Check if a LoRA request can be scheduled.

        This method checks two conditions:
        1. The drainer allows scheduling (based on draining state)
        2. The LoRA adapter can be loaded (either already running or can be added)
        """
        if self.lora_drainer and not self.lora_drainer.can_schedule(req):
            return False

Set --lora-drain-wait-threshold 5.0 and any adapter whose requests have waited more than five seconds triggers a controlled drain of one running adapter to free its slot. Default is 0.0, i.e. disabled.

LoRA and speculative decoding

The pinned SGLang HEAD is literally [Spec][LoRA] Support multi-adapter LoRA with EAGLE/NEXTN/DFLASH/DSPARK speculative decoding (#34337). §6.6 covers that from the draft side; the LoRA side is the constraint stated in the validator:

    def _check_lora_speculative_compatibility(self):
        """Validate LoRA + speculative decoding combinations.

        Adapters apply to the target only; a shared draft runs unadapted.
        Matches resolved algorithm names (NEXTN has collapsed to EAGLE).
        """
        if self.speculative_algorithm in ["NGRAM", None]:
            return

        if self.speculative_algorithm not in _LORA_SPEC_ALGORITHMS:
            promoted = (
                " (NEXTN/EAGLE with a Gemma4 assistant draft is automatically "
                "promoted to FROZEN_KV_MTP, which does not support LoRA)"
                if self.speculative_algorithm == "FROZEN_KV_MTP"
                else ""
            )
            raise ValueError(
                "LoRA is only compatible with NGRAM, EAGLE, NEXTN, EAGLE3, "
                "DFLASH, or DSPARK speculative decoding, not "
                f"{self.speculative_algorithm}{promoted}."
            )

One caveat that reads like a contradiction with §6.6 and is not. DSPARK is the one algorithm whose supports_ragged_verify() returns True (python/sglang/srt/speculative/spec_info.py:L131-L136), and the compact ragged-verify CUDA-graph path refuses to run with LoRA at all: "Compact ragged verify does not support two-batch-overlap, LoRA, or disable-cuda-graph-padding" (python/sglang/srt/model_executor/runner/decode_cuda_graph_runner.py:L350-L358). That is not a blanket LoRA+DSPARK ban, because ragged verify is not on by default: SGLANG_RAGGED_VERIFY_MODE defaults to "static" (python/sglang/srt/environ.py:L1137), and in static mode DSpark's ragged layout is never built at all. Default configuration: LoRA + DSPARK runs.

Be precise about the refusal, though, because the CUDA-graph message above is the narrower of two gates and it is not the one you hit first. The server-arg validator refuses on self.speculative_algorithm == "DSPARK" and ragged_mode != "static" (python/sglang/srt/server_args.py:L9590-L9593) — not equal to "compact". RaggedVerifyMode has three members, STATIC, CAP_ACCEPT and COMPACT (python/sglang/srt/speculative/ragged_verify.py:L13-L16), so cap-accept is refused exactly as compact is. Either one gets you a startup ValueError reading "LoRA with EAGLE/NEXTN/EAGLE3 speculative decoding does not support SGLANG_RAGGED_VERIFY_MODE=..." — note the prefix names EAGLE even when the algorithm is DSpark. The same list refuses --speculative-adaptive under LoRA unconditionally, with no mode condition (server_args.py:L9596-L9601). All of these are startup errors during argument validation, not decode-time corruption.

"Adapters apply to the target only; a shared draft runs unadapted" is the whole design in one sentence, and it has a real consequence: acceptance rate degrades as the adapter moves the target's distribution away from the draft's. A heavily fine-tuned adapter and a base-model draft will accept fewer tokens than the base model would. Nobody has published that curve; it is the obvious experiment.

A batch of three requests: request 0 uses adapter legal (int id 1, 5 tokens scheduled), request 1 uses adapter medical (int id 2, 4 tokens), request 2 uses the base model (3 tokens). Total 12 tokens — the batch in Figure 1.

1. GPUModelRunner.execute_model reaches set_active_loras from the mixin. It asks the input batch to flatten per-request adapter ids into per-token ids:
   vllm/v1/worker/lora_model_runner_mixin.py:L64-L84 vLLM

       def _set_active_loras(
           self,
           prompt_lora_mapping: tuple[int, ...],
           token_lora_mapping: tuple[int, ...],
           lora_requests: set[LoRARequest],
           mapping_type: LoRAMappingType = LoRAMappingType.LANGUAGE,
       ) -> None:
           self._ensure_lora_enabled()
   
           # Set is_prefill to True, so we always use the SGMV kernels on
           # non-cuda platforms.
           # On cuda platforms we use the same kernels for prefill and
           # decode and this flag is generally ignored.
           lora_mapping = LoRAMapping(
               token_lora_mapping,
               prompt_lora_mapping,
               is_prefill=True,
               type=mapping_type,
           )
           self.lora_manager.set_active_adapters(lora_requests, lora_mapping)

   token_lora_mapping is now the 12-element tuple (1,1,1,1,1, 2,2,2,2, 0,0,0) in request order.
2. WorkerLoRAManager.set_active_adapters calls _apply_adapters — the guard from §1 — then add_adapter for anything not yet resident, which loads from disk and calls activate_adapter to claim a slot and copy_ the weights into lora_a_stacked[slot] / lora_b_stacked[slot] for every wrapped module.
3. _set_adapter_mapping calls punica_wrapper.update_metadata(mapping, self.lora_index_to_id, self.lora_slots + 1, self.vocab_size). Note the translation: the tuple holds adapter ids; convert_mapping rewrites them into slot indices using lora_index_to_id, with base-only tokens becoming -1.
4. LoRAKernelMeta.prepare_tensors sorts and uniques (quoted above). Result: lora_ids = [-1, slot(1), slot(2)], num_tokens_per_lora = [3, 5, 4], lora_token_start_loc = [0, 3, 8, 12], token_indices_sorted_by_lora_ids = [9,10,11, 0,1,2,3,4, 5,6,7,8]. There is also a no_lora_flag_cpu fast path: if every token maps to -1, the kernels do not launch at all.
5. Forward runs. At each wrapped layer, _apply_sync does the base GEMM over all 12 rows, then add_lora_linear allocates buffer [n_slices, 12, 16] fp32, launches lora_shrink (grid axis 2 = adapter, so the -1 CTA returns immediately), then lora_expand which accumulates into the base output in place.
6. At the LM head, LogitsProcessorWithLoRA uses sampler_indices rather than token_lora_indices — one entry per sampled position, not per scheduled token, because the head only runs on last positions.

Everything above happens once per forward pass, per adapted module. For Llama-3-8B with all seven targets that is 224 shrink launches and 224 expand launches per step — although the merged QKV and gate/up wrappers fuse their slices into one launch each, bringing it to 128 of each. Kernel-launch overhead is why SGLang went to fixed-size chunks and CUDA-graph capture for the prefill path (PR #30988).

This is the operational fact that surprises people, and it is a direct consequence of the soundness argument in §2.3. Adapted weights produce different K and V for the same token, so KV blocks computed under adapter A are simply wrong for adapter B. The engines say so in one function each:

def _gen_lora_extra_hash_keys(request: Request) -> list[str]:
    """Generate extra keys related to LoRA for block hash computation.

    Args:
        request: The request object.

    Returns:
        Return LoRA name of the request if it is a LoRA request. Return empty
        list otherwise.
    """
    if not request.lora_request:
        return []
    return [request.lora_request.lora_name]

        # Extra key for caller-defined request classification.
        if lora_id is not None:
            extra_key = (
                extra_key or ""
            ) + lora_id  # lora_id is concatenated to the extra key

Quantify it. Take §2.1's budget for the reference machine, and mind the units — an “80 GB” H100 is 79.65 GiB, neither 80 GiB nor 80×10⁹ bytes. At gpu_memory_utilization=0.92 that is 73.28 GiB; subtract 14.96 GiB of Llama-3-8B weights (8.03 B parameters × 2 bytes, converted to GiB rather than left as decimal GB) and 6.0 GiB of activations, graphs and non-torch allocations, and the KV pool is 52.32 GiB. At 128 KiB/token that is 428,569 tokens of cache. A shared 2,048-token system prompt costs 2,048 tokens per tenant, not per deployment:

A miss on the common 2048-token prompt has the stated approximate compute cost, but fifty tenants do not imply misses on every request. Estimate saved work as hit probability times avoided prefill cost, accounting for partial hits and retained last-token computation. Measure reuse distance and evictions rather than inferring an exact TTFT penalty from footprint.

Two mitigations, both real:

Embedding, classification and reranking typically finish after one sequence forward. They avoid the autoregressive decode loop, but still have ragged lengths, queueing, padding, activation limits, cancellation and admission decisions. Short inputs or small batches can be launch- or bandwidth-bound; long bidirectional attention can dominate compute and memory. One-pass completion does not make batching trivial or prove tensor-core saturation.

Pooling strategies, and why the model owns the choice

vLLM's three sequence-pooling methods are exactly the three you would guess:

class CLSPool(SequencePoolingMethod):
    def forward(
        self,
        hidden_states: torch.Tensor,
        pooling_metadata: PoolingMetadata,
    ) -> SequencePoolingMethodOutput:
        pooling_cursor = pooling_metadata.get_pooling_cursor()
        if pooling_cursor.is_partial_prefill():
            raise RuntimeError("partial prefill is not supported with CLS pooling")

        return hidden_states[pooling_cursor.first_token_indices_gpu]


class LastPool(SequencePoolingMethod):
    def forward(
        self,
        hidden_states: torch.Tensor,
        pooling_metadata: PoolingMetadata,
    ) -> SequencePoolingMethodOutput:
        pooling_cursor = pooling_metadata.get_pooling_cursor()
        return hidden_states[pooling_cursor.last_token_indices_gpu]

SGLang's are line-for-line equivalent, differing only in that it computes the indices from extend_seq_lens directly rather than through a cursor object:

    if pooling_type == PoolingType.LAST:
        last_token_indices = torch.cumsum(forward_batch.extend_seq_lens, dim=0) - 1
        return hidden_states[last_token_indices]
    elif pooling_type == PoolingType.CLS:
        prompt_lens = forward_batch.extend_seq_lens
        first_token_flat_indices = torch.zeros_like(prompt_lens)
        first_token_flat_indices[1:] += torch.cumsum(prompt_lens, dim=0)[:-1]
        return hidden_states[first_token_flat_indices]
    elif pooling_type == PoolingType.MEAN:
        prompt_lens = forward_batch.extend_seq_lens
        end_indices = torch.cumsum(prompt_lens, dim=0) - 1
        cumulative_hidden_states = torch.cumsum(hidden_states, dim=0)
        sums = cumulative_hidden_states[end_indices]
        preceding_sums = torch.zeros_like(sums)
        preceding_sums[1:] = cumulative_hidden_states[end_indices[:-1]]
        return (sums - preceding_sums) / prompt_lens.unsqueeze(-1)
    

Pooling should match the checkpoint's training and documented inference contract. BERT-family models are not universally CLS-only: Sentence-BERT variants explicitly train mean-pooled sentence embeddings. An arbitrary replacement can degrade retrieval, but "garbage" is not a theorem. Check special tokens, attention masks, normalization and classification heads, then evaluate recall/ranking metrics. The engine's model defaults and override controls must be interpreted against that model-specific contract.

How you land on the pooling path at all is a separate decision. vLLM's --runner defaults to "auto" and resolves through the model registry (ModelConfig._get_runner_type, vllm/config/model.py:L1169-L1191); a generative checkpoint is turned into an embedder with --convert embed or --convert classify, which forces runner_type = "pooling". SGLang's equivalent is the single boolean --is-embedding, "Whether to use a CausalLM as an embedding model", default False (python/sglang/srt/server_args.py:L600-L602). Neither is a separate runner class: pooling branches inside the same GPUModelRunner / ModelRunner, on the presence of PoolingParams.

What the engine actually has to change

Chunked prefill and prefix caching turn themselves off. The pooling type is not just a correctness property of the model — vLLM reads it back out at config time and uses it to pick the engine's defaults. ModelConfig.is_chunked_prefill_supported and is_prefix_caching_supported (vllm/config/model.py:L2040-L2083, L2086-L2138) return False for any pooling model with attn_type == "encoder_only", and for a causal-attention pooling model whose seq_pooling_type is MEAN or CLS (or whose token pooling is STEP). Those two properties are precisely what _set_default_chunked_prefill_and_prefix_caching_args assigns when the operator has not passed the flags (vllm/engine/arg_utils.py:L2673-L2723). So a BGE-style encoder starts with both features silently disabled, a LAST-pooling decoder-LM embedder starts with both enabled, and forcing the flag the other way logs "This model does not officially support chunked prefill" before letting you do it. This is the mechanism behind the MEAN-pooling RuntimeError in the pitfalls section: the exception is a backstop, not the front line.

No KV reservation for growth. SGLang gates on a single property:

    @property
    def is_prefill_only(self) -> bool:
        """Check if this request is prefill-only (no token generation needed)."""
        # NOTE: when spec is enabled, prefill_only optimizations are disabled

        spec_alg = get_spec().speculative_algorithm
        return self.sampling_params.max_new_tokens == 0 and spec_alg is None

For a genuinely encoder-only model, the saving is total — the KV spec reports zero:

@dataclass(frozen=True)
class EncoderOnlyAttentionSpec(AttentionSpec):
    def max_memory_usage_bytes(self, vllm_config: VllmConfig) -> int:
        # Encoder-only layers do not need KV cache
        return 0

Bidirectional attention. This is the one place the attention backend abstraction of §3.4 has to bend. BERT and RoBERTa are not causal. Both engines carry an AttentionType on the layer and flip the FlashAttention causal flag:

            max_seqlen_k=max_seqlen_k,
            softmax_scale=self.scale,
            causal=False,  # Encoder attention is bidirectional
            alibi_slopes=self.alibi_slopes,
            window_size=sliding_window_size,
            softcap=self.logits_soft_cap,

        causal = not (
            layer.is_cross_attention
            or layer.attn_type
            in (AttentionType.ENCODER_ONLY, AttentionType.DECODER_BIDIRECTIONAL)
        )

and the model declares it at construction:

        self.attn = RadixAttention(
            num_heads=self.num_heads,
            head_dim=self.head_dim,
            scaling=self.scaling,
            num_kv_heads=self.num_kv_heads,
            layer_id=layer_id,
            prefix=f"{prefix}.attn",
            attn_type=AttentionType.ENCODER_ONLY,
        )

Note the second-order effect in SGLang's window-size logic just below that quote: a sliding-window encoder layer needs a symmetric window (w, w), not the causal (w, 0). That kind of detail is why bidirectional attention is not simply "pass causal=False".

Output path returns tensors. vLLM's model runner branches at the top of _pool, and the first thing it does is assert homogeneity:

        num_reqs = self.input_batch.num_reqs
        assert num_reqs == len(self.input_batch.pooling_params), (
            "Either all or none of the requests in a batch must be pooling request"
        )

        hidden_states = hidden_states[:num_scheduled_tokens]
        seq_lens_cpu = self.optimistic_seq_lens_cpu[:num_reqs]

        pooling_metadata = self.input_batch.get_pooling_metadata()
        pooling_metadata.build_pooling_cursor(
            num_scheduled_tokens_np,
            seq_lens_cpu,
            device=hidden_states.device,
            query_start_loc_gpu=self.query_start_loc.gpu[: num_reqs + 1],
        )

        model = cast(VllmModelForPooling, self.model)
        raw_pooler_output: PoolerOutput = model.pooler(
            hidden_states=hidden_states, pooling_metadata=pooling_metadata
        )

Cross-encoder reranking

vLLM's task taxonomy names the three scoring architectures directly:

PoolingTask = Literal[
    "embed",
    "classify",
    "token_embed",
    "token_classify",
    "plugin",
    "embed&token_classify",
]
POOLING_TASKS: tuple[PoolingTask, ...] = get_args(PoolingTask)

ScoreType = Literal["bi-encoder", "cross-encoder", "late-interaction"]
SCORE_TYPE_MAP: dict[PoolingTask, ScoreType] = {
    "embed": "bi-encoder",
    "classify": "cross-encoder",
    "token_embed": "late-interaction",
}

A cross-encoder reranker puts query and document in one sequence — [CLS] query [SEP] document [SEP] — and emits one relevance scalar. That changes the batching shape completely from an embedding server:

• Sequences are short and uniform. A rerank of the top-50 candidates for one query is 50 sequences of maybe 256–512 tokens. Compare that to an embedding workload where a document can be 8k tokens.
• Fan-out is per query, not per request. One user query becomes 50 forward-pass rows. The natural batch is the candidate list, and it arrives as a unit — so a rerank endpoint has a batching opportunity that a generic embedding endpoint does not.
• Bidirectional query representations depend on the document. In a BERT-style cross-encoder, an identical text query does not yield reusable higher-layer query KV across different documents. Causal rerankers have a different dependency structure; inspect the actual attention mask before assuming prefix reuse.

SGLang's cross-encoder pooler is deliberately simple, and it does the per-sequence slicing in a Python loop before a single batched classifier call — see CrossEncodingPooler.forward at python/sglang/srt/layers/pooler.py:L233-L263. Its multi-item scoring path (pool_at_delimiter_positions, python/sglang/srt/layers/pooler.py:L77-L113) goes further: pack many documents into one sequence separated by a delimiter token, then extract hidden states at the pre-delimiter positions and score only those. The comment at the call site says it plainly — "Pool hidden states at pre-delimiter positions, score only those — avoids wasting compute on tokens that never contribute to the output" (score_and_pool, python/sglang/srt/layers/pooler.py:L135-L146). Note the two conditions on that branch: it is taken only when forward_batch.multi_item_delimiter_indices is not None and forward_batch.is_prefill_only, so it is a scoring-endpoint optimisation, not something a generic embedding request gets.

Why these belong in an inference engine at all

Because the machinery is already there and it is 90% of the work. Paged allocation, the tokenizer service, continuous batching, tensor parallelism, CUDA graph capture, the OpenAI-compatible HTTP surface, Prometheus metrics, request cancellation — none of that is generation-specific, and rewriting it for an embedding server is how you end up with two systems that fail differently at 3 a.m. And a RAG deployment needs an embedder, a reranker, and a generator; running all three on the same engine binary means one set of flags, one set of dashboards, and one place to look when latency moves.

Silent: the wrong adapter, applied correctly. The dangerous LoRA bugs never raise. The mapping is a tuple of adapter ids that gets rewritten into slot indices; if slots are reassigned but the mapping tuple compares equal, stale metadata can send a request's tokens through another tenant's weights. vLLM has fixed exactly this (PR #47725, "Refresh punica metadata when LoRA slots are reassigned under an unchanged mapping"). A sibling hazard is documented in the dual-stream path:

        def lora_fn() -> torch.Tensor:
            # Must be zeros, not empty: _lora_expand_kernel exits early (without
            # writing) when lora_id == -1 (no active LoRA). If uninitialized,
            # output.add_(lora_result) below would corrupt the base output.
            lora_output = torch.zeros(
                (num_tokens, output_size),
                device=self.device,
                dtype=x.dtype,
            )

torch.empty instead of torch.zeros there produces garbage logits for base-model requests only, in a code path that only runs when VLLM_LORA_ENABLE_DUAL_STREAM is set. That is a two-day bug.

Silent: int32 overflow at long context. The row-gather index in the shrink kernel is a token index into a flattened [num_tokens, hidden] tensor. Multiply by input_d0_stride = 4096 and a 600k-token batch overflows int32. The fix is the explicit cast you can see in vllm/lora/ops/triton_ops/kernel_utils.py:L296-L297: # int64 to keep the row offset from overflowing at long context lengths (PR #53034). Before it, long-context LoRA silently read from the wrong rows.

Debugging method. For any "adapter seems to be ignored" report: set VLLM_LOGGING_LEVEL=DEBUG and look for Activating LoRA. int id: %d, slot index: %d from model_manager.py:L334-L335. If the slot index churns every step, your working set exceeds --max-loras and you are thrashing the buffers. For "embeddings look wrong", print model.pooler — vLLM's SequencePooler.extra_repr prints pooling=<method>, head=<head>, which tells you in one line whether you are getting CLS when the model wants LAST.

The row-vector formula must agree with merging the stored [out,in] weight. The second half changes an upstream hidden state through an MLP-shaped residual update, while holding the later K projection fixed; the resulting K changes. This is the dependency missed by the claim that adapting only gate/up/down leaves the full KV cache unchanged. For pooling, additionally test masks, empty inputs, left/right padding, normalization and the trained pooling head.

import numpy as np

rng = np.random.default_rng(5)
x = rng.normal(size=(3, 4))
w = rng.normal(size=(6, 4))
a = rng.normal(size=(2, 4))
b = rng.normal(size=(6, 2))
scale = 0.5
unmerged = x @ w.T + scale * (x @ a.T) @ b.T
merged = x @ (w + scale*b@a).T
np.testing.assert_allclose(unmerged, merged, atol=1e-12)
wk = rng.normal(size=(2, 4))
mlp_delta = np.array([0.5, 0., 0., 0.])
base_k = x @ wk.T
adapted_k = (x + mlp_delta) @ wk.T
assert not np.allclose(base_k, adapted_k)
np.testing.assert_allclose(adapted_k-base_k,
                           np.broadcast_to(mlp_delta@wk.T, base_k.shape))
print("Merged LoRA parity passes; upstream MLP changes later keys.")

Serve three adapters on one base model and watch the slot machinery move. No lab exists for this chapter yet; these are the commands.

VLLM_LOGGING_LEVEL=DEBUG vllm serve meta-llama/Meta-Llama-3-8B-Instruct \
  --enable-lora \
  --max-loras 4 \
  --max-lora-rank 16 \
  --max-cpu-loras 16 \
  --enable-prefix-caching \
  --lora-modules legal=/adapters/legal medical=/adapters/medical finance=/adapters/finance

# Run this client in a second terminal after the server is ready:
python - <<'PY'
import json
import urllib.request
from pathlib import Path
prompt = Path("2k_system_prompt.txt").read_text()
for adapter in ("legal", "medical", "finance", "legal"):
    payload = json.dumps({"model": adapter, "prompt": prompt, "max_tokens": 8}).encode()
    req = urllib.request.Request("http://localhost:8000/v1/completions",
        data=payload, headers={"Content-Type": "application/json"})
    with urllib.request.urlopen(req, timeout=120) as response:
        result = json.load(response)
    assert result["choices"]
    print(adapter, result.get("usage"))
PY

The fourth request reuses legal and should show hits; the second and third should show none, despite sending byte-identical prompts. That is Figure 2, measured.

python -m sglang.launch_server --model-path meta-llama/Meta-Llama-3-8B-Instruct \
  --enable-lora \
  --lora-paths legal=/adapters/legal medical=/adapters/medical finance=/adapters/finance \
  --max-loras-per-batch 4 \
  --lora-backend csgmv \
  --max-lora-chunk-size 64 \
  --lora-eviction-policy lru \
  --lora-drain-wait-threshold 5.0 \
  --log-level debug

# compare against the triton backend, which does not chunk:
#   --lora-backend triton

The thing to measure: send a skewed load (90% to legal, 10% split across the others) and compare p99 TTFT between csgmv and triton. The chunking exists specifically for that distribution.

vllm serve BAAI/bge-base-en-v1.5 --runner pooling

curl -s localhost:8000/v1/embeddings -H 'Content-Type: application/json' \
  -d '{"model":"BAAI/bge-base-en-v1.5","input":["a","b","c"]}' | jq '.usage'

# confirm the KV spec really is zero for an encoder-only model:
curl -s localhost:8000/metrics | grep -E 'gpu_cache_usage|kv_cache'

Flag names are volatile — as of a556f3f the runner selector is --runner; check vllm serve --help before assuming.

• Hu et al., LoRA: Low-Rank Adaptation of Large Language Models (arXiv 2106.09685) — the factorisation itself, and the empirical claim that $r$ as low as 1–4 suffices for many tasks.
• Chen et al., Punica: Multi-Tenant LoRA Serving (arXiv 2310.18547) — the SGMV kernel. vLLM cites it by name and arXiv id in the module docstring of vllm/lora/punica_wrapper/punica_base.py:L5-L7, and SGLang's ChunkedSgmvLoRABackend docstring links the PDF directly. The paper reports large throughput gains over per-adapter batching; read its configuration before quoting the number.
• Sheng et al., S-LoRA: Serving Thousands of Concurrent LoRA Adapters (arXiv 2311.03285) — unified paging for adapter weights and KV, and the origin of the GPU-slots-plus-CPU-pool architecture both engines now use.
• SGLang PR #34337 — [Spec][LoRA] Support multi-adapter LoRA with EAGLE/NEXTN/DFLASH/DSPARK speculative decoding. This is the pinned HEAD commit. Read it alongside §6.6.
• SGLang PR #30988 — Support LoRA under the breakable/full prefill CUDA graph, and PR #32708 — Shard attention LoRA by attn-TP and allow dynamic LoRA with DP attention. Together these are why csgmv chunks into fixed sizes.
• vLLM PR #47725 — Refresh punica metadata when LoRA slots are reassigned under an unchanged mapping and PR #53034 — Fix int32 index overflow in LoRA punica kernels at long context. The two canonical silent-wrongness bugs in this subsystem.
• vLLM PR #10507 — Refactor KVCacheManager for more hash input than token ids — where extra_keys was generalised, making the per-adapter cache partition of Figure 2 possible in the first place. Background in §2.3.
• Source to read next: vllm/lora/ops/triton_ops/README_TUNING.md, which documents how the shrink/expand block sizes are autotuned per (max_loras, batch, hidden_size), and python/sglang/kernels/ops/gemm/lora_tuning_config.py for SGLang's equivalent.