Commit Graph

1311 Commits

Author SHA1 Message Date
yyj 52b46826d4 Delete kt-kernel/test/test_swiglu_limit_method_guards.py 2026-08-15 22:00:55 +08:00
yyj 7b1cd66475 feat: support SwiGLU limit for block FP8 experts 2026-08-12 03:13:13 +08:00
github-actions[bot] eb9b70c411 [build]: sync sglang submodule (#2069) 2026-08-08 18:45:46 +08:00
yyj f5ee50d4cf [docs] dsv4 tutorial update(#2140)
* docs: clarify DSV4 Docker runtime defaults
2026-08-05 18:01:01 +08:00
yyj e03cf1cefe docs: add DeepSeek V4 Flash Docker guide (#2136) 2026-08-05 00:11:27 +08:00
yyj 1a14f473cc [build] Sync sglang-kt submodule #2117 2026-08-02 20:11:37 +08:00
yyj a8062bfa7e release: bump version to 0.6.4 (#2105) 2026-07-23 22:42:28 +08:00
yyj 924754a00b [feat](kt-kernel): end-to-end full-parameter and LoRA SFT(#2094)
* Add full FT development snapshot

* [fix](kt-kernel): fix Full FT TP base weight gradients

* [fix]: bug fix of 2d81e86

* [fix](kt-kernel): fix AMX BF16 full-weight gradients

* [docs](kt-kernel): document Full FT fork changes and debug history

* [docs](kt-kernel): align fork remote terminology

* [chore](kt-kernel): keep agent notes local

* [fix](kt-kernel): configure SFT OpenMP threads

* [perf](kt-kernel): optimize AMX Full-FT weight gradients

Coarsen base-weight gradient work from individual output tiles to fixed-intermediate strips so each task reuses packed panels across the hidden dimension.

Keep aligned thread-local BF16 panels across tasks and retain FP32 AMX accumulator tiles for the full K reduction. Gate and up run separate K passes while sharing the packed input panel.

On the matched Qwen3-30B-A3B 1-GPU test, stable Full-FT backward drops from 9.281s to 6.792s (-26.82%), step time drops from 19.252s to 16.140s, and TPS rises from 212.76 to 253.78. The LoRA-only backward control changes by -2.74%.

Validated with clang-format, the Release AMX/CUDA extension build, TP1/TP2 reference gradients across boundary token counts, and the 15-step Full-then-LoRA performance run.

* [feat](kt-kernel): add staged SFT profiling

* [fix](kt-kernel): reuse inference BF16 kernel for SFT

* [perf](kt-kernel): add fine-grained Full-FT profiling

* [perf](kt-kernel): batch BF16 Full-FT weight gradients

Use one expert-aggregated tile driver for AVX512-BF16 and AMX base-weight gradients, and pack updated full-precision weights directly into TP BufferB layouts without temporary partitions. Add worker-local profiling and focused dWeight/strided-repack coverage.

* [test](kt-kernel): benchmark BF16 dWeight AMX driver

* [fix](kt-kernel): label dWeight store as worker CPU time

* [docs](sft): record BF16 Full-FT performance

* [docs](sft): remove Qwen3 Full-FT performance report

* [perf](kt-kernel): reduce BF16 Full-FT checkpoint overhead

Retain the first CPU MoE forward state across non-reentrant checkpoint recomputation, write BF16 activations directly into the backward cache, and reduce dWeight packing and gradient-clear traffic. Extend staged profiling and cover checkpoint reuse plus AMX/AVX dWeight paths.

* [perf](kt-kernel): make SFT optimizer gradients authoritative

Bind Full-FT and LoRA Parameter.grad directly to the KT-managed BF16 gradient buffers, avoiding PyTorch duplicate accumulation.

Accumulate microbatch gradients in C++, lazily clear expert buffers between optimizer windows, preserve rank-0 distributed ownership, and add lifecycle and AMX dWeight coverage.

* [perf](kt-kernel): avoid eager Full-FT gradient zeroing

Allocate authoritative Full-FT gradient buffers with torch.empty. The C++ state machine performs the mandatory full clear before first use, avoiding redundant Python-side first touch.

* perf(sft): enable checkpoint forward reuse for LoRA

* fix(sft): serialize checkpoint recompute with async repack

* fix(sft): persist authoritative full weights

* fix(sft): normalize legacy distributed gradients

* [fix](sft): train gated shared experts

* fix(sft): preserve expert placeholders across state dict loads

* ci: publish ktransformers sdist in release workflow

* fix(sft): preserve router autograd for LoRA training

Determine routing graph tracking from the router's trainable parameters instead of the Full-FT mode flag, so PEFT LoRA adapters on MoE gates receive routing-weight gradients. Fail fast when a trainable router returns detached weights and cover frozen, TopK, checkpoint-reuse, and two-step optimizer behavior.

---------

Co-authored-by: illu <wubowen03@foxmail.com>
v0.6.4
2026-07-23 21:23:34 +08:00
yyj d1a3ed8a30 [perf]: add K2 RAWINT4 prefill mat-mat dispatch (#2080)
* [perf]: add K2 RAWINT4 prefill mat-mat dispatch

* [perf] Refactor K2 RAWINT4 prefill accumulation

---------

Co-authored-by: yyj6666667 <yyj6666667@users.noreply.github.com>
2026-07-19 21:43:01 +08:00
Aaryan Mehta 937f61c988 [fix](kt-kernel): AVX-VNNI-256 RAWINT4 per-expert weight loading (#2092) 2026-07-19 17:24:52 +08:00
Tai An def0f9313d fix(security): bind balance_serve scheduler ZMQ socket to loopback (#2087) (#2091) 2026-07-19 11:38:22 +08:00
callmegaga 0c2912a5e2 [feat] Add SYCL backend for intel iGPU (#2089)
* [feat](sycl): add sycl backend for igpu

* [feat](sycl): improve prefill

* [feat](sycl): finish and test sycl gptq int4
2026-07-17 16:24:04 +08:00
Graham 01fdfa609e docs(DSV4): add supplementary deployment and tuning notes
* docs(DSV4-Flash): AMXINT4 on Skylake (no AMX), dual-Xeon tuning + single-stream negative results

Adds a section based on a validated single-workstation deployment (1x RTX 5090 +
2x Xeon Gold 6138 / Skylake-SP): AMXINT4 runs on AVX512F+BW without an AMX tile;
dual-socket tuning notes; and a table of configs that reduced single-stream
throughput here (notably EAGLE going net-negative on a CPU-bandwidth-bound box,
vs the 8x-RTX-5090 speedup the guide cites).

* Fix command syntax in DeepSeek-V4-Flash.md

---------

Co-authored-by: yyj <159887351+yyj6666667@users.noreply.github.com>
2026-07-15 13:06:58 +08:00
callmegaga 7c021b430c [chore](sglang): bump sglang submodule for FP16 FE8M0 Marlin fix (#2082) 2026-07-10 11:18:45 +08:00
VectorPeak cb9f47d142 [fix](cli): detect bound ports before launch (#2071)
* [fix](cli): detect bound ports before launch

* [fix](cli): align port reuse check by platform
2026-07-06 18:25:31 +08:00
lutianshu824 79b265b2f6 fix: normalize compressed RAWINT4 weights (#2075)
* fix: normalize compressed RAWINT4 weights

* docs: add Hygon DCU ROCm notes

---------

Co-authored-by: lutianshu824 <lutianshu824@users.noreply.github.com>
2026-07-06 18:06:52 +08:00
Benjamin 8e46e5896c release: bump version to 0.6.3.post1 (#2063) v0.6.3.post1 2026-06-25 22:42:55 +08:00
github-actions[bot] 459df61262 [build] Sync sglang-kt submodule (v0.6.3) (#2055) 2026-06-25 20:41:04 +08:00
Benjamin a0c7431187 ci(release-pypi): make release pipeline self-consistent (#2062) 2026-06-25 20:05:06 +08:00
Willow Lopez 983a88b620 docs: add FAQ entry about SGLang MoE kernel config warning (#2029) 2026-06-24 15:12:54 +08:00
Hermann Hans Klie 08d70e8605 [fix](kt-kernel): enable -mavx512vl for AVX512_VBMI/BF16 multi-variant builds (#2021) 2026-06-24 15:05:36 +08:00
Jiaheng Dai 1ed332b6ea [docs]: clean up Qwen3.5 KT LoRA serving guide (#2057)
* [docs]: clean up Qwen3.5 KT LoRA serving guide

* Update doc/zh/Qwen3.5-SGLang-LoRA-Serving_zh.md

Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>

---------

Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
2026-06-23 20:52:23 +08:00
Benjamin c884dbd221 release: bump version to 0.6.3 (#2054) 2026-06-21 17:55:24 +08:00
github-actions[bot] ce7c3ddbe9 [build]: sync sglang submodule to 8b636f9008dbad58c0a8e481b03e794739e6c146 (#2047)
Co-authored-by: ovowei <80044717+ovowei@users.noreply.github.com>
v0.6.3
2026-06-21 17:37:52 +08:00
Benjamin 6c9c95601d [fix](kt-kernel): CudaGraph replay fix and Add MoE startup log (#2037)
CudaGraph replay fix and Add MoE startup log
2026-06-18 17:15:37 +08:00
Benjamin 943cc4daeb [feat] MXFP8 MoE support (#2041)
Add MXFP8 kernel for Minimax-M3-MXFP8 Day 0
2026-06-18 16:18:49 +08:00
Tai An 21d5b40323 docs(deepseek-v4-flash): pin apache-tvm-ffi<0.1.12 to avoid tilelang TVM FFI registration clash (#2045) (#2048) 2026-06-18 11:10:56 +08:00
Benjamin 50d9434800 [docs](kt-kernel): add MiniMax-M3 SGLang + KT-Kernel tutorial (#2051) 2026-06-18 10:40:51 +08:00
Jianwei Dong 0f2e7905e2 Update README.md (#2049) 2026-06-17 14:24:23 +08:00
Jianwei Dong 0e2ff783d8 update glm52 tutorial (#2046) 2026-06-16 17:15:23 +08:00
Dayuxiaoshui 7641f5445d [feat] : Add MACA backend support for kt-kernel and fix CPU MoE tests (#2044)
* Add MACA backend support for kt-kernel

* Add MACA event API mappings

* Fix AMX build flags and CPU MoE tests


---------

Co-authored-by: <Engle_Chaveztih@sociologist.com>
2026-06-16 16:49:06 +08:00
devangpratap 89d30a3d01 [fix(loader)]: correct off-by-one expert-count guard in load_experts (#2026)
* [fix(loader)]: correct off-by-one expert-count guard in SafeTensorLoader.load_experts

After the discovery loop, max_experts_count is the highest expert index found
(expert count - 1), and is -1 only when the key has no experts. The guard
checked == 0, which falsely rejected single-expert layers and silently returned
empty weight lists for the zero-expert case. Check == -1 instead.

Adds a CPU regression test covering the single-, zero-, and multi-expert cases.

* [test(loader)]: import loader as a top-level module in expert-count guard test

Per review feedback: add python/utils to sys.path and import loader directly
instead of the importlib.util boilerplate. Still bypasses utils/__init__.py
(and the compiled kt_kernel_ext) while keeping the import idiomatic.
2026-06-07 23:41:04 +08:00
github-actions[bot] c1cb22311b [build]: sync sglang submodule to 51032b71279d9038058563f8d2e758d99b278ef4 (#2032)
Co-authored-by: ovowei <80044717+ovowei@users.noreply.github.com>
2026-06-05 19:24:22 +08:00
Jiaheng Dai c9a915e6ac [feat](kt-lora): add end-to-end Qwen3.5 MoE KT LoRA serving workflow (#2031)
* [feat](kt-lora): add KT expert LoRA adapter serving

* [feat]: pin Qwen3.5 non-expert LoRA support

* [feat](kt-lora): add merged SGLang adapter workflow

Document the KT SFT to SGLang serving loop and extend the converter with optional split outputs so users can serve one merged adapter while retaining debug-friendly expert/non-expert artifacts.

Co-authored-by: Cursor <cursoragent@cursor.com>

* [fix](kt-lora): validate adapter conversion

Co-authored-by: Cursor <cursoragent@cursor.com>

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-06-05 16:57:14 +08:00
devangpratap d41f569e84 [fix](cli): detect SGLANG_DSV4_2604_SUBMODE conflict before launch (#2025)
* [fix](cli): detect SGLANG_DSV4_2604_SUBMODE conflict before launch

* [fix](cli): tighten env-var validation per review feedback

doctor.py: skip SGLANG_DSV4_2604_SUBMODE row when value is empty string,
not just None, to avoid spurious noise in kt doctor output.

run.py: guard kt_method against None/empty before calling .upper() in
_check_conflicting_env_vars to prevent AttributeError.
2026-05-30 19:20:47 +08:00
Benjamin ef6c47f9d2 [feat](kt-kernel): AVX2 MXFP4 MoE MXFP4 dispatch (#2015)
* [feat](kt-kernel): AVX2 MXFP4 MoE MXFP4 dispatch

- Add AVX2 MXFP4 MoE kernel (mxfp4-moe.hpp) with 4-token M-blocking
- Wire AVX2MXFP4_MOE binding in ext_bindings.cpp
- Support TP_MOE down_proj slicing and multi-pool per-expert loading
- Add test_fp4_moe_avx2.py integration test

* [fix](kt-kernel): address PR #2010 review — memory leaks, alignment, dynamic expert update

- Track aligned_alloc pointers in AVX2_MOE_BASE::owned_aligned_allocs_ and
  free them in the destructor (fixes BufferB backing memory leak on destroy).
- Track per-TP down_buf allocations in TP_MOE::tp_owned_down_bufs_ with
  nullptr checks and size rounding to alignment boundary.
- Add nibble-alignment runtime check for per_tp_interm in MXFP4 TP K-split.
- Add write_weight_scale_to_buffer override to TP_MOE<AVX2_MXFP4_MOE_TP>,
  enabling dynamic expert update with kt-threadpool-count>=2.
- Guard against ZeroDivisionError in test_fp4_moe_avx2.py.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

* [fix](kt-kernel): add intermediate_size parity check in MXFP4 TP flat-buffer path

The per-expert path validates that intermediate_size is even (required for
nibble-aligned FP4 addressing), but the flat-buffer path was missing this
check — an odd value would silently truncate /2 divisions, corrupting
memcpy sizes and offsets.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

* fix(avx2-moe): fix TP offset calculation and add safety checks

C1-C4: Fix incorrect TP offset calculations in load_weights()
- Per-expert mode used per_tp_interm instead of full_interm for offsets
- This caused segfault when TP > 1 due to invalid pointer arithmetic

H1-H3: Add safety checks
- H1: Validate source weight pointers are not null
- H2: Check lid index is within bounds
- H3: Check BufferB.b is not null in gemm_mxfp4

Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

* fix(avx2-moe): revert incorrect C2/C4 offset changes, keep safety checks

Reverts the incorrect offset calculation changes from previous commit.
The original per_tp_interm-based offsets were correct:
- gate/up weights are N-split (along intermediate dim)
- Each TP partition handles per_tp_interm rows
- Offset = i * per_tp_interm * hidden / 2 (not full_interm)

Keeps H1-H3 safety checks:
- H1: Validate source weight pointers are not null
- H2: Check lid index is within bounds
- H3: Check BufferB.b is not null in gemm_mxfp4

Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

* fix(avx2): copy weights to owned buffers in per-expert mode

Previously, AVX2 MXFP4 MoE per-expert mode directly pointed BufferB.b
into mmap'd safetensor data. This caused use-after-free when Python
layer releases the mmap after load_weights() returns.

Now AVX2 copies weights into owned buffers via memcpy/from_raw_mat(),
matching AMX behavior. This decouples the MoE weights from mmap lifecycle.

Changes:
- buffer_b_required_size_impl: always allocate full buffer (weights + scales)
- make_buffer_b_impl: always create full BufferB with owned storage
- Single-TP per-expert: use from_raw_mat() instead of direct pointer
- TP_MOE per-expert: add gate/up owned buffers with memcpy
- Destructor: free gate/up buffers alongside down

Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

* Revert "[fix] Add runtime AMX BF16 check to prevent SIGILL on pre-Sapphire Rapids CPUs (#2018)"

This reverts commit f1e2b82c74.

* Remove AMX tile MXFP4 kernel (GemmKernel224MXFP4)

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

---------

Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
2026-05-30 19:20:16 +08:00
Li Tingfang f1e2b82c74 [fix] Add runtime AMX BF16 check to prevent SIGILL on pre-Sapphire Rapids CPUs (#2018) 2026-05-21 17:36:12 +08:00
login256 eeeeae5e91 Fix duplicate BF16 loader definition (#1984) 2026-05-20 15:04:47 +08:00
Jim James f0772445a1 [perf]: native path for MXFP4 MoE on AVX512F (#2006)
* [perf]: native path for MXFP4 MoE on AVX512F

* [perf]: move inline static constants outside structs
2026-05-18 15:44:33 +08:00
github-actions[bot] 95e20f9c55 [build]: sync sglang submodule to ebaff7729b9e41c29d94f8d19a53473d321dc566 (#2005)
Co-authored-by: ovowei <80044717+ovowei@users.noreply.github.com>
2026-05-14 22:25:31 +08:00
Benjamin F f05b4009f3 [fix](kt-kernel): fix double mem used by safetensor loader (#1997)
Release the SafeTensor mmap loader singleton after each layer's
load_weights() completes. The C++ engine already holds a deep copy
(cpu_infer.sync() guarantees this), so releasing the mmap handles is
safe. The next layer recreates the loader on demand.

This halves peak memory usage during model loading (e.g. DSv3.2:
1.2T -> 613G).

Based on #1966 by @poryfly — adapted to v0.6.2.post3 codebase
(adds MXFP4 support missing from the original PR).

Co-authored-by: xiongchenhui <xiongchenhui@hisense.com>
2026-05-11 12:00:30 +08:00
Benjamin F bb15fdf47e Release/0.6.2.post3: carry kt-kernel SwiGLU clamp companion missing from post2 v0.6.2.post3 2026-05-10 03:55:02 +08:00
Benjamin F 37db9a3b83 0.6.2.post2: submodule refactor and update tutorial (#1993)
- sglang submodule -> 43ed1ec77: V4-Flash hybrid SWA chunked-prefill
  hang fix (#44) + DSV4 plugin registry refactor (#47)
- pin sglang-kt==0.6.2.post2
- tutorial: switch V4-Flash launch example from 8x RTX 5090 to
  single-card (decode 20+ tok/s); flip Ada Lovelace SM_89 row to
  validated; update Hardware Requirements GPU line accordingly
v0.6.2.post2
2026-05-09 18:53:59 +08:00
Jim James f7c4fa68c5 [fix]: add guard for SFT MoE and remove guard for AMX FP4 MoE on AVX512F+BW (#1980) 2026-05-08 16:05:22 +08:00
Benjamin F c465557c23 docs(v4-flash): add optional AMXINT4 CPU-weight conversion path (#1986)
- Add convert_cpu_weights_ds4.py: dequantizes MXFP4 routed experts
  (E2M1 + ue8m0, group size 32) on GPU and re-quantizes to AMX-INT4 on CPU.
- Document the script as Step 2 in DeepSeek-V4-Flash.md so AMX users
  can opt into AMXINT4 mode instead of the default MXFP4 CPU experts.
2026-05-08 15:35:05 +08:00
Benjamin F 8b9d233d42 docs(v4-flash): tilelang install, MTP flags, Ampere unsupported (#1979)
- Prerequisites: add `pip install tilelang` (sglang-kt extras don't
  pull it; required for the NSA tilelang indexer fallback on
  non-Hopper GPUs). Validated with tilelang==0.1.8.
- Step 2: add an "Optional: Enable MTP" subsection with the EAGLE
  speculative-decoding flags for V4-Flash NextN.
2026-05-06 17:29:38 +08:00
Benjamin F d7b5b49a3e [release]: 0.6.2.post1
V4-Flash MXFP4 full-GPU prefill fallback now works:
- Previously crashed all TP schedulers with StopIteration/AttributeError
  whenever --kt-gpu-prefill-token-threshold was low enough to actually
  fire (path was hardcoded for FP8/INT4 layouts).
- Now detects MXFP4, re-runs the V4 swizzle on the 256-expert gpu_layer,
  caches the load across prefill chunks.
- Measured on 8x RTX 5090 (threshold=1024, chunked=1024):
  16k input -> 2011 tok/s, 65k -> 2798, 262k -> 2154 prefill TPS.
v0.6.2.post1
2026-05-03 21:07:23 +08:00
Benjamin F 96189972d8 build: bump sglang submodule to c9edb75e0 (V4-Flash GPU prefill fallback fix + perf) (#1975) 2026-05-03 19:42:19 +08:00
Benjamin F 088ed979d5 docs(v4-flash): pin transformers==4.57.1 in tutorial prerequisites (#1974)
V4-Flash is incompatible with the transformers 5.x series. transformers
5.x adds default-valued fields to PretrainedConfig that make
DeepSeekV4Config's dataclass declaration crash at import time with:

    TypeError: non-default argument 'quantization_config' follows default argument

Reproduced on a fresh venv with `pip install sglang-kt`: pip resolves
transformers to 5.7.0 (sglang-kt's pyproject does not pin transformers),
and `python -m sglang.launch_server --model .../DeepSeek-V4-Flash` fails
during import of sglang.srt.configs.deepseek_v4. Pinning to 4.57.1 fixes
it.

Add a Prerequisite #4 documenting the explicit pin alongside the existing
flashinfer override.
2026-05-03 16:07:31 +08:00
Benjamin F 4b4312c0a2 release: bump version to 0.6.2 (#1973) v0.6.2 2026-05-03 14:28:09 +08:00