SGLang 源码完全拆解（十三）：多 LoRA 服务——adapter 池、SGMV 内核与批内多适配器共存

三层映射

1
lora uid  <-->  buffer_id (0 .. max_loras_per_batch-1)  <-->  物理权重缓冲

三个结构维持这个映射：

• uid_to_buffer_id: Dict[Optional[str], int]：adapter 名字到槽号；
• buffer_id_to_uid: List[Union[str, None, EmptySlot]]：槽号到 adapter 名字，空槽用 EmptySlot 单例占位；
• A_buffer / B_buffer: Dict[str, List[torch.Tensor]]：按目标模块组织的物理张量。

注意 buffer_id_to_uid 里同时存在三种值：str 是真实 adapter，None 是基座模型（无适配器），EmptySlot 是空槽。源码注释专门解释了为什么用 EmptySlot 而不是 None：

All uids are initialized as EmptySlot for empty buffer slots. Here we don’t initialize to None since None is a valid uid.

这是一个很实际的细节。None 在业务上表示「这个请求用基座模型」，所以它必须能占用槽位；如果空槽也用 None 表示，就无法区分「槽里装的是基座」和「槽是空的」。用一个独立的哨兵对象解决歧义。

物理缓冲的形状

A_buffer / B_buffer 是 Dict[str, List[torch.Tensor]]，外层键是目标模块名（如 q_proj、down_proj），值是一个列表——列表的第 i 个元素就是第 i 号槽的缓冲区。所以每个模块的所有 adapter 权重挤在一段连续显存里，索引方式就是 A_buffer["q_proj"][buffer_id]。

单个缓冲区的形状取决于模块类型：

1
标准 LoRA（3D）: [num_loras, rank, hidden_dim]
2
MoE LoRA（4D）:  [num_loras, num_experts, rank, hidden_dim]

源码注释直接给出了这两种形状。MoE 多一维是因为 FusedMoE 的权重本来就是按专家分组的，LoRA 修正必须逐专家施加。

还有一个容易被忽略的点：rank 维度按 max_lora_rank 分配，而不是每个 adapter 自己的 rank。源码里 load_lora_weight_tensor 会先断言 buffer_view.shape == weight.shape 再拷贝，而截断逻辑在内核侧由 N = tl.minimum(N, rank * stack_num) 完成。这意味着不同 rank 的 adapter 共用同名槽形状，rank 小的 adapter 右侧留空。这是以空间换通用性：不为每个 adapter 单独分配精确大小的缓冲，从而让槽位可以互换、内核模板可以复用。

并行维度的分片口径

缓冲区形状必须精确匹配基座权重在那一个 rank 上的分片。这里有三套容易混淆的并行度：

源码里有一段很长的注释专门讲这件事，值得完整理解：

Per-expert MoE weights are sharded by moe_tp_size, NOT the outer tp_size: moe_tp_size = tp_size // ep_size // moe_dp_size, so under e.g. --tp 4 --ep 4 each rank holds full-width expert weights (moe_tp_size == 1).

也就是说，在 --tp 4 --ep 4 的配置下，每个 rank 持有的是完整宽度的专家权重（因为专家已经被 EP 切开了，不再需要 TP 再切），所以 LoRA buffer 的内维也应该是完整宽度。如果这里用 tp_size 去算，会得到窄四倍的内维，与 adapter 权重的形状对不上，加载时直接形状断言失败。

这类「分片口径」问题是多 LoRA 实现里最容易出错的地方，因为它不报语义错误，只报形状错误，而且只在特定并行组合下出现。

初始化与容量校验

can_support() 在初始化时校验 adapter 是否放得下，validate_lora_batch() 在运行时校验当前批次是否需要超过 max_loras_per_batch 个 adapter。如果超出，SGLang 的选择是直接抛错而不是静默降级：

1
The global DP-attention batch contains more LoRA adapters than the memory pool can hold

这条约束在 DP attention 场景下尤其重要。因为 DP attention 下每个 rank 只看到自己那一份请求，但最终要 all-gather 出一个全局批次，所以校验必须在 gather 之后、全局视图上进行——源码中 active_lora_ids 正是通过 tp_group.all_gather_object(forward_batch.lora_ids) 汇总出来的：

1
if self.attn_dp_enabled:
2
    gathered_lora_ids = get_parallel().tp_group.all_gather_object(forward_batch.lora_ids)
3
    active_lora_ids = {
4
        lora_id
5
        for rank_lora_ids in gathered_lora_ids
6
        for lora_id in rank_lora_ids
7
    }
8
    if not self.validate_lora_batch(active_lora_ids):
9
        raise ValueError(...)
10
    self.fetch_new_loras(active_lora_ids)

注意这里是按集合汇总的。同一时刻全局有多少个不同的 adapter，而不是有多少个请求——因为一个 adapter 可以被多个请求共用，槽位只需要一个。

从朴素公式出发

对某层的一个线性模块，设输入激活为 $X \in \mathbb{R}^{s \times k}$（$s$ 是批次里所有 token 拼起来的行数），基座权重为 $W \in \mathbb{R}^{d \times k}$（PyTorch 习惯，nn.Linear 里 $W$ 的形状是 [out_features, in_features]），LoRA 权重为 $A \in \mathbb{R}^{r \times k}$、$B \in \mathbb{R}^{d \times r}$。单个适配器的输出是：

其中 $\alpha$ 是 LoRA 的缩放系数（通常 $= \text{lora\_alpha} / r$，SGLang 在 LoRAConfig 里读取 lora_alpha 并在加载时算出 scaling）。

现在假设批次里有 $n$ 个不同的适配器，token 按适配器分成了 $n$ 个连续的段（segment）。段 $i$ 有 $s_i$ 行，$\sum_i s_i = s$。分段矩阵向量乘（SGMV，有时也叫 SGMM 分组矩阵乘，因为这里实际处理的是矩阵而非向量）要算的是：

下标 $i$ 表示「用第 $i$ 个适配器的权重」。基座项对所有段都一样，是一个普通的 GEMM；LoRA 项是 $n$ 段各用一套不同 $A_i, B_i$ 的分块 GEMM。

为什么分两步：shrink 与 expand

直接算 $(XA_i^\top)B_i^\top$ 当然可以，但有个问题：$A$ 和 $B$ 的形状不同，最优的 tile 划分也不同。

• $XA^\top$：形状 $s_i \times r$，是「压缩」——把 $k$ 维输入降到 $r$ 维。r 通常只有 16 或 32，很小。这个 GEMM 的 K 维很大（k=4096），M 维是 token 数，N 维很小。
• $(XA^\top)B^\top$：形状 $s_i \times d$，是「展开」——把 $r$ 维升回 $d$ 维。这个 GEMM 的 K 维很小（r=16），M 维是 token 数，N 维很大。

两个 GEMM 的瓶颈完全不同。前者受 K 维循环的算术强度支配，后者受 N 维输出的写带宽支配。Punica 的核心洞察就是：把它们当作两个独立的内核来调度，中间结果放在显存里做中转。于是有了 shrink kernel 和 expand kernel。

这与 vLLM 源码完全拆解（八）：MoE 层的前向执行 里 grouped GEMM 的拆解思路是同源的——把「每个分组用不同权重」的困难，转化为「一段连续输出里按索引取不同权重」的困难，然后让内核自己去做这个索引。

shrink kernel：_sgemm_lora_a_kernel

SGLang 的 Triton 实现把这一步写成：

1
@triton.jit
2
def _sgemm_lora_a_kernel(
3
    x, weights, output,
4
    N,        # stack_num * r
5
    K,        # input_dim
6
    stack_num,
7
    x_stride_0, x_stride_1,
8
    w_stride_0, w_stride_1, w_stride_2,
9
    output_stride_0, output_stride_1,
10
    seg_lens, seg_indptr, weight_indices, lora_ranks, sorted_token_ids,
11
    SORTED_BY_ADAPTER: tl.constexpr,
12
    BLOCK_S: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr,
13
):

逐个参数看过来，几个关键点：

1. batch_id = tl.program_id(axis=1) —— 每个 program（即一次网格调度）负责一个段。这是 SGMV 与普通 GEMM 在内核组织上的根本区别：网格的第二维是「段」，而不是「批」。
2. w_index = tl.load(weight_indices + batch_id) —— 用段号查表得到适配器槽号。这一条 load 就是「批内多适配器」的全部魔法：内核不需要知道有哪些 adapter，只需要一张 seg -> buffer_id 的查找表。
3. rank = tl.load(lora_ranks + w_index) —— 再查一次得到该适配器的实际 rank。
4. if rank == 0: return —— rank 为 0 意味着这个请求用基座模型，LoRA 分支的输出应为零，直接跳过整个 program。这是一个干净的处理方式：空操作比特殊分支更符合内核的执行模型。
5. N = tl.minimum(N, rank * stack_num) —— 把输出宽度裁到实际 rank。stack_num 是「一次处理多少个适配器分片」的堆叠因子，见下文的 stack 讨论。

内核主体的循环就是标准的分块 GEMM：

1
    s_offset = tl.arange(0, BLOCK_S) + pid_s * BLOCK_S
2
    n_offset = tl.arange(0, BLOCK_N) + pid_n * BLOCK_N
3
    k_offset = tl.arange(0, BLOCK_K)
4
    x_ptrs = x + (s_physical[:, None] * x_stride_0 + k_offset[None, :] * x_stride_1)
5
    w_ptrs = (weights + w_index * w_stride_0) + (
6
        k_offset[:, None] * w_stride_2 + n_offset[None, :] * w_stride_1
7
    )

注意 x_ptrs 用的是 s_physical 而不是 s_offset。s_physical 由 _resolve_token_positions() 计算，它把逻辑行号映射到物理行号——这就是 permutation 的作用。如果批次里的 token 顺序不是按适配器整理好的，内核需要先做一次重排。

w_ptrs 的基址是 weights + w_index * w_stride_0，即先跳到该适配器的权重块，再按 $(k, n)$ 索引。这里的 weights 就是前面 A_buffer[module_name] 那个列表在显存里拼成的二维视图。

最后：

1
    partial_sum += tl.dot(x_tile, w_tile)

tl.dot 用 FP32 累加。这一点和 FlashAttention 完全拆解（一） 里讨论过的在线 softmax 数值考虑一致：低精度输入、高精度累加是这类内核的通用做法。

expand kernel 与累加

expand kernel（_sgemm_lora_b_kernel）把 shrink 的输出 $\tilde{X} = XA^\top \in \mathbb{R}^{s \times r}$ 与 $B$ 相乘，得到最终修正 $\tilde{X}B^\top \in \mathbb{R}^{s \times d}$。它的索引逻辑与 shrink 相同（同样按段查表），区别在 tile 划分：K 维是 $r$ 而不是 $k$，N 维是 $d$ 而不是 $r$，所以循环次数与 block 配置不一样。源码里的 csgmv_configs 目录按 lora_shrink / lora_expand 分别存放调优好的配置：

1
lora_expand,K=1024,R=64,S=1,device=NVIDIA_H200.json
2
lora_shrink,K=3072,R=64,S=1,device=NVIDIA_H200.json

文件名里的 K、R、S 就是内核的三个形状参数（K 维、rank、stack_num），是按实际负载形态预调好的 autotune 结果。

图源：Punica 论文 Figure 3（arXiv:2310.18547）。图中把「一组请求的输入」与「一组适配器权重」按段对齐，输出则是同一个 GEMM 的形状。段的边界由 seg_indptr 给出，段的权重由 weight_indices 给出——这两个张量就是 SGMV 接口的全部输入。

累加落点：为什么修正项要「加」上去

还有一个需要交代的设计点：SGMV 内核的输出是修正项，不是最终结果。基座 GEMM 仍由普通路径算（SGLang 里就是标准 Linear 前向），LoRA 修正以 += 的形式加在结果上。

这带来一个重要的性能含义：如果整个批次都不用 LoRA，那么完全不进入 LoRA 路径。LoRAManager.prepare_lora_batch() 开头就有这个快速返回：

1
if self.lora_backend.skip_inactive_lora_batches and not any(
2
    uid is not None for uid in forward_batch.lora_ids
3
):
4
    self.reset_lora_batch()
5
    return

skip_inactive_lora_batches 是后端的一个属性，注释说明它是给 UNO 这类「用显式 token 行路由」的内部后端用的。对 Triton 后端而言，即使不跳过，rank == 0 的早退也会让所有段立刻返回，代价只是网格启动本身。

这也解释了 SGLang 文档里那一句提醒：「enabling LoRA on additional modules introduces a minor performance overhead」——每多一个目标模块，就多两次内核启动（shrink + expand）和一段缓冲区，即使批内没有任何请求使用它。

seg_indptr：段边界

标准的 CSR（Compressed Sparse Row）式前缀和表示。段 $i$ 覆盖的行区间是 [seg_indptr[i], seg_indptr[i+1])。这个设计的好处是：段长可以不等——prefill 批次里不同请求的 prompt 长度天差地别，用 indptr 表示比用固定的 seg_len 数组更自然。注意源码注释指明 Triton 后端下 num_segments 恰好等于 bs，即「一个请求一段」。

但 chunked backend（CSGMV）下两者会不相等，源码注释解释了原因：

for csgmv they differ because its segments are chunked across adapters.

也就是说 CSGMV 会把一个段再切块，用更细的粒度换取更好的负载均衡。这就是为什么除了 seg_indptr 之外还有一套 req_seg_indptr：

1
    # Per-request segment indptrs, shape (bs + 1,). Required by MoE virtual
2
    # experts which map tokens to requests regardless of the dense-LoRA
3
    # backend's internal segmentation.
4
    req_seg_indptr: Optional[torch.Tensor] = None
5
    req_weight_indices: Optional[torch.Tensor] = None

两套索引的存在原因：不同消费者需要不同的分组视图。dense LoRA 内核按「后端自己的分段」工作（可能与请求边界不一致），而 MoE 的虚拟专家机制必须按「请求边界」工作。与其让某一方将就，不如把两套都算出来。

permutation：逻辑行到物理行

这是最容易被误解的字段。它回答的问题是：内核看到的输入行号，与输出张量里的物理行号，是不是同一个顺序？

如果内核把段按请求顺序处理，而输入激活是按调度顺序排列的，两者可能不一致。此时 permutation[i] 给出「第 i 个逻辑行对应哪个物理行」。shrink kernel 里那一行：

1
    s_physical = _resolve_token_positions(
2
        sorted_token_ids, seg_start, s_offset, seg_len, SORTED_BY_ADAPTER
3
    )

就是 permutation 的应用点。当 SORTED_BY_ADAPTER 为真时，token 已经按适配器排好序，位置解析是恒等映射，_resolve_token_positions 直接返回 s_offset；否则需要真查表。

这是一个重要的性能开关：让 token 按适配器排序可以把 gather 变成顺序读，代价是一次重排。SGLang 把它做成 constexpr，说明这是编译期决定的策略，而不是运行时分支。

lora_ranks 与 scalings：按槽索引，不是按段

注意这两个张量的形状是 (lora_num,) 而不是 (num_segments,)——它们按槽号索引：

1
    lora_ranks = [0] * self.max_loras_per_batch
2
    scalings = [0] * self.max_loras_per_batch
3
    for uid in active_lora_ids:
4
        if uid is None:
5
            continue
6
        weight_index = self.memory_pool.get_buffer_id(uid)
7
        lora = self.loras[uid]
8
        lora_ranks[weight_index] = lora.config.r
9
        scalings[weight_index] = lora.scaling

weight_indices 把段映射到槽，lora_ranks / scalings 把槽映射到属性。两次间接寻址，换来的是内核侧只读两张小表，而不是遍历请求列表。

scalings 的默认值是 0——这意味着如果某个槽没有被正确赋值，它的缩放系数是 0，修正项会算出 0 而不是脏值。这是一个保守的初始化选择：宁可少算，不要算错。

weight_indices 的默认值：基座槽

1
    base_weight_index = self.memory_pool.uid_to_buffer_id.get(None, 0)
2
    weight_indices = [base_weight_index] * len(forward_batch.lora_ids)
3
    for i, uid in enumerate(forward_batch.lora_ids):
4
        if uid not in self.memory_pool.uid_to_buffer_id:
5
            continue
6
        weight_indices[i] = self.memory_pool.get_buffer_id(uid)

整个数组先填成「基座模型所在槽」，再逐个替换成实际槽号。注意那个 continue：如果某个 uid 不在池里（理论上不该发生，因为前面已经 fetch 过），它保持基座槽。

uid_to_buffer_id.get(None, 0) 里的默认值 0 也值得注意——None 在初始化时就被放进了映射（占一个槽），所以正常情况下一定命中；默认值只是防御性写法。

has_active_lora：给 Python 侧的一个布尔开关

1
    self.lora_backend.batch_info.has_active_lora = any(
2
        lora_ranks[wi] > 0 for wi in weight_indices
3
    )

注释说明它是「computed from Python lists in prepare_lora_batch to avoid GPU sync」。这是一个纯 CPU 侧的计算，因为 lora_ranks 在构造时还是 Python list。这样 Python 层的后续逻辑（比如要不要给 lm_head 加 LoRA、要不要走特殊分支）可以直接读这个布尔值，不需要把张量同步回 CPU。在推理服务里，任何一次隐式的 .item() 都是一次同步点，代价是整个流水线停一下——所以「用 Python list 算好再转张量」是这类系统的常规手法。

第一段：加载（磁盘 → CPU）

load_lora_adapter() 的日志很能说明问题——它在开始和结束各打一次可用显存：

1
logger.info(
2
    f"LoRA adapter loading starts: {lora_ref}. "
3
    f"avail mem={get_available_gpu_memory(self.device.type, self.device.index):.2f} GB"
4
)

在加载前后记录显存，是为了让运维能立刻看到「加载一个 adapter 到底吃了多少显存」。这类日志在多 LoRA 服务里非常必要，因为 adapter 的内存占用与 max_loras_per_batch、max_lora_rank、目标模块集合三者相乘，不实际测一下很难估准。

加载分两步：先读配置（LoRAConfig），再读权重（load_lora_weights）。配置里最重要的四个字段：

1
self.target_modules = self.hf_config["target_modules"]  # 适配哪些模块
2
self.r = self.hf_config["r"]                            # 秩
3
self.lora_alpha = self.hf_config["lora_alpha"]           # 缩放分子
4
self.use_dora = self.hf_config.get("use_dora", False)    # 是否 DoRA

target_modules 决定这个 adapter 需要往哪些缓冲区写数据；r 和 lora_alpha 决定缩放。

LoRAConfig 里还有一处不显眼但很实用的处理——过滤伪 added tokens：

1
    # Filter fake added tokens: tokens with ID < base_vocab_size are already
2
    # part of the base vocabulary and should not be treated as added tokens.
3
    # This commonly happens when added_tokens.json is copied from the base
4
    # model's tokenizer.
5
    if self.added_tokens_config and base_vocab_size is not None:
6
        self.added_tokens_config = {
7
            token: token_id
8
            for token, token_id in self.added_tokens_config.items()
9
            if token_id >= base_vocab_size
10
        }

有些 adapter 的发布包里带了一份从基座模型复制来的 added_tokens.json，里面的 token 全部属于基座词表。如果不加过滤，系统会以为这个 adapter 有自己的新增 token 表，从而在 embedding 层分配多余的缓冲区。

validate_new_adapter() 还会做兼容性校验。源码里有一条与 DP attention 相关的硬约束：

1
assert not self.attn_dp_enabled or lora_ref.pinned, "DP-attention LoRA requires pinned adapters"

DP attention 下必须 pin 住 adapter。原因不难推：DP attention 下每个 rank 只持有自己那部分请求，但全局批次需要所有 rank 对「有哪些 adapter」达成一致；如果 adapter 可以被某个 rank 单方面淘汰，不同 rank 的全局视图就会不一致，all_gather_object 之后的校验也就失去意义。pinned 是一个跨 rank 的一致性保证。

第二段：换入（CPU → GPU 槽位）

fetch_new_loras() 是换入的入口：

1
    def fetch_new_loras(
2
        self, new_loras: set[Optional[str]], running_loras: set[Optional[str]] = set()
3
    ):
4
        cur_uids = new_loras | running_loras
5
        assert len(cur_uids) <= self.max_loras_per_batch
6
        new_uids = {
7
            uid for uid in cur_uids if uid not in self.memory_pool.uid_to_buffer_id
8
        }
9
        self.memory_pool.prepare_lora_batch(
10
            cur_uids=cur_uids,
11
            lora_adapters=self.loras,
12
            lora_modules=self.lora_modules,
13
            lora_refs=self.lora_refs.copy(),
14
            lora_embed_tokens_module=self.embed_tokens_module,
15
            lora_lm_head_module=self.lm_head_module,
16
        )

三个要点：

1. cur_uids = new_loras | running_loras —— 换入的候选集合包含「本批次新出现的」加上「当前正在跑的」。语义是：正在跑的 adapter 不能被换出，否则它的后续 decode 步骤会读到错误的权重。
2. new_uids 用差集算出 —— 只有不在池里的才需要真正换入。
3. lora_refs.copy() —— 源码注释解释了原因：copy snapshot of current lora_refs to avoid mutation during the batch preparation。prepare_lora_batch 内部会修改池状态，传一份快照避免遍历时被并发修改。这是一个典型的「先快照再修改」防御。

换入之后还要通知各层模块：

1
    def _notify_lora_slots_updated(self, slot_ids: set[int]) -> None:
2
        for layer_modules in self.lora_modules:
3
            for module in layer_modules.values():
4
                notify = getattr(module, "on_lora_slots_updated", None)
5
                if callable(notify):
6
                    notify(slot_ids)

为什么需要通知？因为有些后端会缓存「槽位到权重张量」的映射（比如 CUDA Graph 捕获时把 buffer 的地址固化进了图里）。槽位内容变了，缓存必须失效。用 getattr(..., None) + callable 的鸭子类型检查，是为了兼容不带这个钩子的模块实现。

第三段：淘汰

当槽位不够时，必须有 adapter 被换出。策略在 eviction_policy.py 里，是一个标准的策略模式：

1
class EvictionPolicy(ABC):
2
    @abstractmethod
3
    def mark_used(self, uid: Optional[str]) -> None: ...
4
    @abstractmethod
5
    def select_victim(self, candidates: Set[Optional[str]]) -> Optional[str]: ...
6
    @abstractmethod
7
    def remove(self, uid: Optional[str]) -> None: ...

两个实现：LRUEvictionPolicy 用 OrderedDict 记录访问顺序，FIFOEvictionPolicy 记录插入顺序。select_victim 从 candidates 里挑一个最老的。注意它的兜底逻辑：

1
        # If no tracked UID found in candidates, check if None is available
2
        # This happens when the batch consists entirely of LoRA requests
3
        # and None (base model) is the only eviction candidate
4
        if None in candidates:
5
            return None

如果候选里只有基座槽，就淘汰基座槽。这对应一个具体的场景：整个批次全是 LoRA 请求，没有任何请求用基座模型，那么基座槽就是唯一可以被腾出来的。此时把基座权重从缓冲区清掉（_clear_buffer_slot_for_base），给新 adapter 让位。

load_lora_weight_tensor 里对应地做了零初始化：

1
            if weight is None:
2
                # If the particular weight is not present in the adapter, we initialize the buffer to zero
3
                # to avoid contamination from the residual weight of the evicted adapters.
4
                buffer_view.zero_()

为什么必须清零？因为槽位是复用的。如果新 adapter 没有某个模块的权重（比如它只适配注意力层，不适配 MLP），那块缓冲区里还留着上一个被淘汰的 adapter 的数据。如果不清零，内核会把这个残留当作新 adapter 的修正项加进去——结果是一个静默的、看不出来的错误输出。这是复用缓冲区最危险的失效模式，源码注释专门点了出来。

权重装载的完整路径

load_lora_weight_to_buffer() 是装载的主函数，它做三件事：

（1）预校验权重名与目标模块的匹配

1
        skipped_weight_names: set = set()
2
        matched_modules: set = set()
3
        ...
4
        for name in all_weight_names:
5
            try:
6
                target_module = get_target_module_name(name, self.target_modules)
7
                matched_modules.add(target_module)
8
            except ValueError:
9
                skipped_weight_names.add(name)

源码注释说明了动机：This catches mismatches before any GPU buffers are mutated. 先全量校验、再写入 GPU，避免中途失败留下半写状态。strict_loading 为真时直接抛错，否则只记警告——两种模式对应「生产要严格」和「兼容性优先」两种诉求。

（2）按层组织临时缓冲

1
            temp_A_buffer: Dict[str, Union[torch.Tensor, Dict[int, torch.Tensor]]] = {
2
                target_module: None for target_module in self.A_buffer
3
            }

这里的值类型是 Union[torch.Tensor, Dict[int, torch.Tensor]]，因为 MoE 权重是「模块 → 专家 → 张量」的三层结构，而普通模块是「模块 → 张量」两层。源码用 re.search(r"experts\.(\d+)\.", name) 从权重名里抽出专家号：

1
                expert_match = re.search(r"experts\.(\d+)\.", name)
2
                is_shared_expert = "shared_experts." in name
3
                shared_moe_target = f"{target_module}_shared_moe"

共享专家（shared expert）被单独命名成 {module}_shared_moe，与路由专家（{module}_moe）分开存放。这与前面拆过的 DeepSeekMoE 完全拆解 里「共享专家恒被激活、路由专家按 Top-K 选择」的结构对应：共享专家的 LoRA 修正对每个 token 都要施加，路由专家只对分到它的 token 施加，两者的索引范围不同，放在同一个缓冲区反而不好索引。

（3）拷入并处理 rank 差异

前面的 copy_weight_into_buffer(buffer_view, weight) 是实际的拷贝，utils.py 里定义。它要处理的核心情形是：adapter 的 rank 小于 max_lora_rank，或者权重的堆叠因子与缓冲区不一致。

第一处：TP——权重的分片口径

这一点在前面「并行维度的分片口径」已经展开。这里补充 S-LoRA 论文给出的通用分片规则。

图源：S-LoRA 论文 Figure 4（arXiv:2311.03285）。核心规则是：

• 对列并行层（q_proj、k_proj、v_proj、gate_proj、up_proj，输出维被切分），$A$ 需要在输出维上切分，$B$ 保持完整——因为 $A$ 的输出维就是 $r$，与权重切分无关；而 $B$ 的输出维是 $d$，必须跟着切。
• 对行并行层（o_proj、down_proj，输入维被切分），情况相反。

SGLang 里对应的实现是 _column_parallel_lora_b_per_rank_dim() 和 get_lora_B_shape() 中的分支。理解这条规则有一个统一的口径：LoRA 修正必须与基座权重在同一 rank 上具有相同的输入/输出维度语义，否则 $W + BA$ 的加法就对不上。

第二处：前缀缓存——adapter 是命名空间的一部分

这一处最容易被忽略，后果却很严重。看 RadixKey 的字段：

1
    __slots__ = ("token_ids", "extra_key", "cache_salt", "is_bigram", "limit")

以及匹配时的断言：

1
        if self.extra_key != other.extra_key:
2
            raise ... f"RadixKey operations require matching extra_key"

源码注释里写得更明确：

extra_key values are intentionally kept disjoint and never share

extra_key 是哈希键的一部分：

1
        return plain if self.extra_key is None else (self.extra_key, plain)

为什么 LoRA 必须进前缀缓存的键？ 考虑一个具体场景：两个请求的 prompt 完全相同，但一个用 adapter A、一个用 adapter B。如果前缀缓存只看 token id，第二个请求会命中第一个请求留下的 KV，于是它读到的是 adapter A 算出来的 KV，却在 adapter B 的权重下继续生成。输出会出错，而且错得很隐蔽——因为 KV 看起来「是合法的」。

所以 adapter 身份必须进键。这也意味着：用 adapter A 和 adapter B 处理同一段 prompt，KV 不能复用。这是多 LoRA 服务的一个真实成本：适配器越多、切换越频繁，前缀缓存的收益越低。反过来，如果 workload 是「同一批用户反复用同一个 adapter 处理相似文档」，那么复用率很高。

这也解释了 lora_name 出现在 kv_event_lora_names 里——SGLang 源码完全拆解（五） 拆过的 KV 事件记账，需要知道每个 KV 页属于哪个 adapter。

第三处：CUDA Graph——形状与地址都必须是静态的

LoRABatchInfo.use_cuda_graph 这个字段的存在就说明问题。CUDA Graph 的核心约束是捕获时的一切都固定：指针地址、形状、分支走向。三个具体的冲突点：

（1）槽位地址

如果换入 adapter 时重新分配缓冲区，地址就变了，图失效。所以 A_buffer / B_buffer 在初始化时就按 max_loras_per_batch 一次性分配好，整个生命周期地址不动。换入只是往固定地址写数据。这正是 _notify_lora_slots_updated 存在的意义——地址没变但要通知缓存失效。

（2）批次形状

use_cuda_graph 的分支要在捕获前决定：

1
    def _use_cuda_graph_batch(self, forward_batch: ForwardBatch) -> bool:
2
        return (
3
            hasattr(self, "max_bs_in_cuda_graph")
4
            and forward_batch.batch_size <= self.max_bs_in_cuda_graph
5
            and forward_batch.forward_mode.is_cuda_graph()
6
            and (not self.attn_dp_enabled or forward_batch.can_run_decode_cuda_graph)
7
        )

四个条件同时满足才能走图。其中 max_bs_in_cuda_graph 是预先探测的批次上限——超过就退回 eager。

（3）原地更新而不是重建

注释里写得很清楚：

1
        # Do in-place updates when CUDA graph is enabled and the batch forward mode
2
        # could use CUDA graph.

以及

1
        # Eligible extend batches refresh the static prefill batch info in
2
        # place so captured kernels read current values at replay.

「原地刷新静态 batch info」 是关键技术：图的拓扑固定，但填入图里的元数据可以在重放前刷新。这样同一个捕获的图可以服务不同的批次内容——只要形状一致。

第四处：MoE——虚拟专家

前面提到 req_seg_indptr / req_weight_indices 是专为 MoE 准备的。这里展开为什么需要。

MoE 层的计算不是「一个 token 一份权重」，而是「一个 token 选 K 个专家，每个专家一份权重」。如果这个 MoE 层又装了 LoRA，就变成「token × 专家 的每一对，都要施加对应 adapter 在这个专家上的修正」。

SGLang 的解法叫虚拟专家：把「(adapter, expert)」的组合当作一个虚拟的专家来索引。源码注释解释了它与 dense LoRA 分段的差异来源：

1
    # For the triton backend these are identical to seg_indptr/weight_indices;
2
    # for csgmv they differ because its segments are chunked across adapters.

对 Triton 后端，虚拟专家的索引恰好等于普通分段（一个请求一段）；对 CSGMV 后端，因为分段被切得更细，虚拟专家必须回溯到请求粒度。同一个语义，两种实现，所以需要两套索引。

这里还牵出一个并行度细节：moe_use_local_expert_ids。

1
        self.moe_use_local_expert_ids = (
2
            self.moe_ep_size > 1
3
            and not _moe_runner_keeps_global_expert_ids()
4
            and num_experts_global % self.moe_ep_size == 0
5
        )

注释说明了原因：StandardDispatcher 会在 MoE 内核之前把全局 topk_ids 重映射成本地专家 ID，所以逐专家的 LoRA 缓冲区必须按本地分片来分配和索引。而 FlashInfer 的 CUTLASS / CuteDSL / TRTLLM / MXFP4 路径保持全局 ID，因此必须按全局索引。同一个 EP 配置，不同的 MoE runner，需要不同的专家编号口径——这是读 SGLang MoE 相关源码时最容易踩的坑之一，可参考 vLLM 源码完全拆解（八）：MoE 层的前向执行 里对 grouped_topk 与专家编号的拆解对照阅读。

第五处：调度——adapter 饥饿与 drain

多 LoRA 服务有一个特有的调度问题：如果 max_loras_per_batch = 2，而队列里有 10 个不同 adapter 的请求，就只能同时跑 2 个。更糟的情况是：少数几个热门 adapter 的长请求一直占着槽位不放，其他 adapter 的请求永远排不上。

SGLang 的解法是 LoRADrainer，源码里 AdapterStats 记录了每个 adapter 的等待情况：

1
class AdapterStats:
2
    def is_starving(self, drain_wait_threshold: float):
3
        ...

对应 --lora-drain-wait-threshold 参数（文档原文）：

When any LoRA adapter request waits longer than this threshold (in seconds), the scheduler will selectively drain one running adapter to make room. This mitigates extreme tail latency under high or skewed workloads by preventing a small set of adapters from monopolizing batch slots.

_update_draining_loras() 里有一条判断：

1
        running_adapter_ids = {req.lora_id for req in running_reqs}
2
        if len(running_adapter_ids) < self.max_loras_per_batch:

只在运行中的 adapter 数少于槽位数时才考虑 drain——因为这时有槽位可以腾出来给饥饿的 adapter。如果槽位已经满载且全部在用，drain 也只是换一个受害者，解决不了根本矛盾。

对应地，_update_fully_drained_loras() 在 adapter 的所有请求都退出时释放槽位。注意这里用 running_reqs 而不是「已经完成的请求」，因为一个 adapter 可能在多个运行中请求里被共用；只有它最后一个请求退出，槽位才真正空闲。

这也和前面的 fetch_new_loras 形成闭环：cur_uids = new_loras | running_loras 里的 running_loras 就是「当前还有请求在跑的 adapter」，它们不会被当作换出候选。

图源：S-LoRA 论文 Figure 3（arXiv:2311.03285）。统一内存池把三样东西放在同一段显存里统一管理：基座权重、adapter 权重（分页存放）、KV 缓存。SGLang 的实现里 adapter 权重所在的 A_buffer / B_buffer 就是图中 adapter 池的对应物，只是它的分页单位是「槽」而不是「页」。

关键参数

陷阱一：max_lora_rank 事后无法放大

源码注释点出了这个约束：

This argument is needed when you expect to dynamically load adapters of larger LoRA rank after server startup.

max_lora_rank 在初始化时就决定了缓冲区形状。如果启动后想加载一个 rank 比它大的 adapter，加载会失败。所以要按「未来可能出现的最大 rank」来设定，而不是按当前已知的 adapter。

同理适用于 lora_target_modules：

This argument is needed when you expect to dynamically load adapters of different target modules after server startup.

如果启动时只声明了 q_proj，之后想加载一个也适配 MLP 的 adapter，那些 MLP 权重会被跳过（skipped_weight_names 逻辑）——adapter 会「加载成功」但行为是错的（MLP 部分没生效）。这是一个特别隐蔽的问题，源码里那段警告就为此而生：

1
                f"LoRA adapter '{uid}': {len(skipped_weight_names)} weight(s) "
2
                f"skipped because they did not match any target module in "

实践建议：如果不确定未来会加载什么 adapter，宁可启动时声明的目标模块集合大一些，承担一点性能开销，也不要让后续加载的 adapter 静默失效。

陷阱二：DP attention 下 adapter 必须 pinned

前面提到 assert not self.attn_dp_enabled or lora_ref.pinned。这条约束的用户侧含义是：在 DP attention 部署下，--lora-paths 里的 adapter 必须用 JSON 形式带 "pinned": true。文档给的是这个格式：

1
{"lora_name":str, "lora_path":str, "pinned":bool}

如果沿用了简化格式（<NAME>=<PATH>），在 DP attention 下会因为 pinned=None 而启动失败。

陷阱三：adapter 隔离带来的缓存收益下降

前面讲过 adapter 进前缀缓存键的必然性。工程上需要注意的推论是：多 adapter 混跑会稀释前缀缓存的收益。如果业务允许把同一 adapter 的请求路由到同一实例，前缀命中率会显著提高。这与 SGLang 源码完全拆解（五）：KV 内存池与 RadixCache 的所有权闭环 里讨论过的前缀匹配逻辑直接相关——命中率取决于有多少请求真正共享 (adapter, tokens) 这一对。

陷阱四：加载时要预留显存余量

回看那段日志——avail mem 打印在加载前后。实践意义是：--max-loras-per-batch × max_lora_rank × target_modules 决定了需要预留的显存。如果启动时把 KV 缓存池按满显存分配，加载新 adapter 时会 OOM。SGLang 在启动时就按 max_loras_per_batch 预分配了缓冲（而不是按需分配），所以这部分显存在启动时就被锁定——不会在运行中突然不够。