Skip to content

fix(peft): correct dense MultiLoRALinear on replicated base linears - #25

Merged
yushengsu-thu merged 1 commit into
bridgefrom
radixark/bridge/fix-multilora-replicated-base
Jul 26, 2026
Merged

fix(peft): correct dense MultiLoRALinear on replicated base linears#25
yushengsu-thu merged 1 commit into
bridgefrom
radixark/bridge/fix-multilora-replicated-base

Conversation

@yushengsu-thu

Copy link
Copy Markdown
Collaborator

Why

A GLM-5.2_5layer multi-LoRA e2e smoke (the first MLA model through the multi-LoRA path) crashed in the first training forward. The dense MultiLoRALinear mishandles replicated base linears — TELinear parallel_mode='duplicated', i.e. the MLA q/kv down-projections of GLM-5 / DeepSeek-style models. Qwen3-style models never reach these paths (all their LoRA targets are TP-parallel linears), which is why the existing e2e coverage missed them.

This was found on #23's branch and pushed there as 592f5e27, but #23 was merged at f15719c1 without it — this PR is that same commit cherry-picked onto the current bridge tip.

Two defects, both specific to replicated bases

  1. Per-slot token spans covered the full micro-batch, but a replicated base under sequence parallelism consumes the SP shard directly (TE stores parallel_mode=None for duplicated linears, so the adapter attribute derivation disables the SP gather). The grouped GEMM's offsets then exceed the row count — an out-of-bounds read surfacing as CUDA error: device-side assert triggered, not a shape error. The spans are now intersected with this rank's contiguous token window (_narrow_token_counts_to_window, the same invariant the MoE slot routing uses for its SP narrow). The host-side total is cached once per micro-batch in set_tokens_per_adapter_slot, so the forward needs no device sync.

  2. _gather_output keyed only on input_is_parallel, leaving the adapter output as a [tokens, out/tp] shard while a replicated base produces the full [tokens, out]. It now mirrors ParallelLinearAdapter's lin_out_gather_output: gather for row-parallel and replicated bases.

Deliberately unchanged: the TP collective between the two grouped GEMMs stays unconditional. ParallelLinearAdapter shards A on the rank axis whenever input_is_parallel=False — replicated bases included — so the mid all-gather is required there too (gating it on base_linear_is_parallel was tried and breaks GEMM2 with a contraction mismatch).

Testing

  • 4 new CPU tests for the window-narrowing math (slot-boundary-spanning windows, identity, single-slot, many-small-slots); the three multi-LoRA test files pass 74/74 on an H200 devbox.
  • GLM-5.2_5layer multi-LoRA e2e (TP=EP=4, ETP=1, SP, thd/tilelang DSA, dp-attention rollout, adapter rank 16 < slot max 32): with this fix (plus [sglang-miles] Allow load_lora_adapter_from_distributed under dp attention sgl-project/sglang#32421 on the serving side) the run completes 3 training steps end-to-end — per-step distributed adapter upserts to both dp-attention engines, (tp, pp, ep) checkpoint shards at the realized coordinates — executed under CUDA_LAUNCH_BLOCKING=1 so no async assert hides.
  • Without the fix, the identical run dies in the first wrapped MLA down-projection (torch._grouped_mm device-side assert), reproduced and localized with CUDA_LAUNCH_BLOCKING=1.

Found by a GLM-5.2 multi-LoRA e2e smoke (MLA q/kv down-projections are
TELinear parallel_mode='duplicated'); Qwen3-style models never exercise
this case because all their LoRA targets are TP-parallel linears.

Two defects, both specific to replicated bases:

* The per-slot token spans covered the full micro-batch, but a replicated
  base under sequence parallelism consumes the SP shard directly (no
  gather), so the grouped GEMM read out of bounds — a device-side assert,
  not a shape error. The spans are now intersected with this rank's
  contiguous token window (the same invariant the MoE slot routing uses
  for its SP narrow), with the host-side total cached once per micro-batch
  in set_tokens_per_adapter_slot so the forward needs no device sync.

* _gather_output keyed only on input_is_parallel, leaving the adapter
  output as a [tokens, out/tp] shard while a replicated base produces the
  full [tokens, out]. It now mirrors ParallelLinearAdapter's
  lin_out_gather_output: gather for row-parallel AND replicated bases.

The TP collective between the two grouped GEMMs stays unconditional:
ParallelLinearAdapter shards A on the rank axis whenever
input_is_parallel=False, replicated bases included, so the mid all-gather
is required there too.

Verified: 74 unit tests pass (4 new for the window narrowing), and a
GLM-5.2_5layer multi-LoRA e2e (TP=EP=4, ETP=1, SP, thd, dp-attention
rollout) completes 3 training steps with per-step adapter upserts and
(tp, pp, ep) checkpoint shards, run with CUDA_LAUNCH_BLOCKING=1.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Copilot AI review requested due to automatic review settings July 26, 2026 04:40

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants