[PR #99] [MERGED] Support April V4 LoRA on original FP8/MXFP4 weights #99

New issue

Closed

opened 2026-10-29 21:07:11 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-29 21:07:11 +00:00

Owner

📋 Pull Request Information

Original PR: https://github.com/kvcache-ai/sglang/pull/99
Author: @yyj6666667
Created: 9/28/2026
Status: ✅ Merged
Merged: 9/30/2026
Merged by: @yyj6666667

Base: main ← Head: yyj/dsv4-flash-delivery

📝 Commits (10+)

e7ca3bf feat(kt): serve static DeepSeek V4 MXFP4 LoRA adapters
948776c fix(kt): retain CPU expert outputs on the CUDA consumer stream
f2ee355 test(attention): cover nonfinite invalid V4 compressed KV slots
7c01294 build: pin private V4 delivery candidate dependencies
fe968db fix: avoid unused TileLang dependency in V4 Triton RoPE
1d746a0 fix(packaging): declare the V4 compressed indexer TileLang dependency
70a62d7 fix(packaging): pin the compatible TileLang TVM FFI pair
228cc9d Format the new static adapter validation and compressed KV test
ad9a88b Load April V4 training adapters on the original quantized base
f675949 Align native adapter tests with lint conventions

📊 Changes

11 files changed (+688 additions, -30 deletions)

View changed files

📝 python/pyproject.toml (+5 -4)
📝 python/sglang/srt/configs/deepseek_v4.py (+2 -1)
📝 python/sglang/srt/layers/deepseek_v4_rope.py (+0 -8)
📝 python/sglang/srt/layers/moe/kt_ep_wrapper.py (+52 -4)
➕ python/sglang/srt/lora/kt_dsv4.py (+218 -0)
📝 python/sglang/srt/lora/lora_manager.py (+13 -0)
📝 python/sglang/srt/models/deepseek_v4.py (+38 -0)
📝 python/sglang/srt/server_args.py (+62 -12)
📝 test/registered/attention/test_dsv4_flash_mla_sm120.py (+31 -1)
➕ test/unit/test_dsv4_kt_static_lora.py (+78 -0)
➕ test/unit/test_kt_dsv4_checkpoint.py (+189 -0)

📄 Description

Serve April DeepSeek-V4-Flash training checkpoints on the original FP8 / MXFP4 base with --lora-paths trained=<checkpoint>. Users no longer need a BF16 export or the training cache for this path. The loader verifies checkpoint hashes, base identity and tensor shapes, maps seven ordinary LoRA targets and reads fused CPU expert LoRA directly. The existing BF16 export path remains available.

The direct path supports one static r8 / alpha16 / dropout0 adapter. It omits only self_attn.compressor.kv_proj and gate_proj LoRA, without skip logs or checkpoint changes; other unsupported tensors fail validation. Requests must select the paired adapter. Base identity currently requires the original training model files. Normal serving without an adapter retains its existing weight-loading and forward paths.

Earlier changes in this PR provide BF16 non-expert loading, clamp-aware MXFP4 SFT dispatch, CPU-output CUDA stream lifetime handling, and the V4 indexer dependencies (tilelang==0.1.10, apache-tvm-ffi==0.1.11). The direct-checkpoint addition changes five runtime files and adds one unit-test file; it requires no native-kernel, KT, Transformers or LF changes.

Validation:

61 adapter-loading/static-request contract tests pass. Changed runtime files pass unused-import/undefined-name checks; the two new files pass the repository lint rules.
Fresh CPython 3.11/3.12 serving environments install with normal hash-locked resolution and pass pip check. Both full training locks also resolve in dry runs. Only the SGLang Python carrier changes; its native CUDA partition and the other 258 wheels are byte-identical to the previous candidate.
All ten direct-loading configurations pass: CPython 3.11/3.12 × AVX512-BF16 1/2/4/8 RTX 5090 GPUs and forced-AVX2 2 GPUs. Each covers single/multi-turn chat, exact 1920+128 and 2047+1 token boundaries, finite logprobs, actual CPU-extension selection and rejection of incomplete base requests. A 10-step training smoke produced the checkpoint with falling loss; its hashes remain unchanged, and the original training-cache path was absent during serving.
No-adapter old/new wheel comparison: six generation fixtures have identical text, token IDs and recorded logprobs; multi-turn chat is identical. Warmed S512/D64 measurements show no regression (six samples per wheel); logged weight/pool memory snapshots match. This does not establish a causal speedup or peak-memory equivalence.
Expert-only and ordinary-only controls both affect logits. Separate BF16 references isolate the original-quantized/runtime-LoRA path difference from the deliberate compressor omission. Neither bitwise equivalence nor quality preservation is claimed. The legacy BF16 expert loader also passed live serving probes.

The previous candidate passed ten training/resume/BF16-serving configurations and Kimi RAWINT4 serving. This update repeats the direct-loading inference matrix, not all training regressions. Forced AVX2 uses an AVX512-capable host; this is not AVX2-only hardware certification. Short training and small probes are functional checks, not general answer-quality acceptance.

Draft; do not merge or publish. Draft GPU CI and existing all-files lint limitations remain. Tested wheel source: ad9a88bc7; the subsequent formatting commits through 1d6f5c6ee preserve the runtime AST. The latest change fixes only PR-added wrapping/import grouping reported by CI; it does not reformat unrelated baseline files.

Related: kvcache-ai/ktransformers#2216 (follow-up to #2185), kvcache-ai/transformers#8. This remains the only SGLang PR for the delivery; no LF PR is requested.

_{🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.}

## 📋 Pull Request Information **Original PR:** https://github.com/kvcache-ai/sglang/pull/99 **Author:** [@yyj6666667](https://github.com/yyj6666667) **Created:** 9/28/2026 **Status:** ✅ Merged **Merged:** 9/30/2026 **Merged by:** [@yyj6666667](https://github.com/yyj6666667) **Base:** `main` ← **Head:** `yyj/dsv4-flash-delivery` --- ### 📝 Commits (10+) - [`e7ca3bf`](https://github.com/kvcache-ai/sglang/commit/e7ca3bf4fdb57ee8c5bbf93cd01c000ff325f4ff) feat(kt): serve static DeepSeek V4 MXFP4 LoRA adapters - [`948776c`](https://github.com/kvcache-ai/sglang/commit/948776cb33a8d1cab68ee993453ccaa51bc12eab) fix(kt): retain CPU expert outputs on the CUDA consumer stream - [`f2ee355`](https://github.com/kvcache-ai/sglang/commit/f2ee35559fd247f30acdbbe52a3552e7c5aa70bf) test(attention): cover nonfinite invalid V4 compressed KV slots - [`7c01294`](https://github.com/kvcache-ai/sglang/commit/7c0129442b5936e08676dc894828b022bb8c8ed6) build: pin private V4 delivery candidate dependencies - [`fe968db`](https://github.com/kvcache-ai/sglang/commit/fe968db401fccba8ca01b9fc0337133c18b89484) fix: avoid unused TileLang dependency in V4 Triton RoPE - [`1d746a0`](https://github.com/kvcache-ai/sglang/commit/1d746a068e99578bc20ccac6c89669aecdf6c0fc) fix(packaging): declare the V4 compressed indexer TileLang dependency - [`70a62d7`](https://github.com/kvcache-ai/sglang/commit/70a62d78928beba6aff78b969c932695216c95ae) fix(packaging): pin the compatible TileLang TVM FFI pair - [`228cc9d`](https://github.com/kvcache-ai/sglang/commit/228cc9dfd35c7227df53148c31f427f681a657b4) Format the new static adapter validation and compressed KV test - [`ad9a88b`](https://github.com/kvcache-ai/sglang/commit/ad9a88bc701d72851cb793c330573472f3f43ee9) Load April V4 training adapters on the original quantized base - [`f675949`](https://github.com/kvcache-ai/sglang/commit/f675949841dcb89789a075c3d928f7392dac33ba) Align native adapter tests with lint conventions ### 📊 Changes **11 files changed** (+688 additions, -30 deletions) <details> <summary>View changed files</summary> 📝 `python/pyproject.toml` (+5 -4) 📝 `python/sglang/srt/configs/deepseek_v4.py` (+2 -1) 📝 `python/sglang/srt/layers/deepseek_v4_rope.py` (+0 -8) 📝 `python/sglang/srt/layers/moe/kt_ep_wrapper.py` (+52 -4) ➕ `python/sglang/srt/lora/kt_dsv4.py` (+218 -0) 📝 `python/sglang/srt/lora/lora_manager.py` (+13 -0) 📝 `python/sglang/srt/models/deepseek_v4.py` (+38 -0) 📝 `python/sglang/srt/server_args.py` (+62 -12) 📝 `test/registered/attention/test_dsv4_flash_mla_sm120.py` (+31 -1) ➕ `test/unit/test_dsv4_kt_static_lora.py` (+78 -0) ➕ `test/unit/test_kt_dsv4_checkpoint.py` (+189 -0) </details> ### 📄 Description Serve April DeepSeek-V4-Flash training checkpoints on the original FP8 / MXFP4 base with `--lora-paths trained=<checkpoint>`. Users no longer need a BF16 export or the training cache for this path. The loader verifies checkpoint hashes, base identity and tensor shapes, maps seven ordinary LoRA targets and reads fused CPU expert LoRA directly. The existing BF16 export path remains available. The direct path supports one static r8 / alpha16 / dropout0 adapter. It omits only `self_attn.compressor.kv_proj` and `gate_proj` LoRA, without skip logs or checkpoint changes; other unsupported tensors fail validation. Requests must select the paired adapter. Base identity currently requires the original training model files. Normal serving without an adapter retains its existing weight-loading and forward paths. Earlier changes in this PR provide BF16 non-expert loading, clamp-aware MXFP4 SFT dispatch, CPU-output CUDA stream lifetime handling, and the V4 indexer dependencies (`tilelang==0.1.10`, `apache-tvm-ffi==0.1.11`). The direct-checkpoint addition changes five runtime files and adds one unit-test file; it requires no native-kernel, KT, Transformers or LF changes. Validation: - 61 adapter-loading/static-request contract tests pass. Changed runtime files pass unused-import/undefined-name checks; the two new files pass the repository lint rules. - Fresh CPython 3.11/3.12 serving environments install with normal hash-locked resolution and pass `pip check`. Both full training locks also resolve in dry runs. Only the SGLang Python carrier changes; its native CUDA partition and the other 258 wheels are byte-identical to the previous candidate. - All ten direct-loading configurations pass: CPython 3.11/3.12 × AVX512-BF16 1/2/4/8 RTX 5090 GPUs and forced-AVX2 2 GPUs. Each covers single/multi-turn chat, exact 1920+128 and 2047+1 token boundaries, finite logprobs, actual CPU-extension selection and rejection of incomplete base requests. A 10-step training smoke produced the checkpoint with falling loss; its hashes remain unchanged, and the original training-cache path was absent during serving. - No-adapter old/new wheel comparison: six generation fixtures have identical text, token IDs and recorded logprobs; multi-turn chat is identical. Warmed S512/D64 measurements show no regression (six samples per wheel); logged weight/pool memory snapshots match. This does not establish a causal speedup or peak-memory equivalence. - Expert-only and ordinary-only controls both affect logits. Separate BF16 references isolate the original-quantized/runtime-LoRA path difference from the deliberate compressor omission. Neither bitwise equivalence nor quality preservation is claimed. The legacy BF16 expert loader also passed live serving probes. The previous candidate passed ten training/resume/BF16-serving configurations and Kimi RAWINT4 serving. This update repeats the direct-loading inference matrix, not all training regressions. Forced AVX2 uses an AVX512-capable host; this is not AVX2-only hardware certification. Short training and small probes are functional checks, not general answer-quality acceptance. **Draft; do not merge or publish.** Draft GPU CI and existing all-files lint limitations remain. Tested wheel source: `ad9a88bc7`; the subsequent formatting commits through `1d6f5c6ee` preserve the runtime AST. The latest change fixes only PR-added wrapping/import grouping reported by CI; it does not reformat unrelated baseline files. Related: kvcache-ai/ktransformers#2216 (follow-up to #2185), kvcache-ai/transformers#8. This remains the only SGLang PR for the delivery; no LF PR is requested. --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>

bsakaguchi

2026-10-29 21:07:11 +00:00

closed this issue
added the
pull-request
label

No labels

pull-request

No milestone

No project

No assignees

1 participant

Notifications

Due date

The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference

kvcache-ai/sglang#99

No description provided.

Rows
Columns