* feat(sft): add native MXFP4 routed-expert LoRA kernels
* test(sft): validate DeepSeek V4 MXFP4 routed experts
* feat(sft): AVX2 tier for MXFP4 routed-expert LoRA SFT
MXFP4_SFT_MOE on non-AVX512 builds, same Python contract as the AMX /
AVX512-BF16 kernel from #2185 (KTMoEWrapper(method="MXFP4_SFT")), for
AVX2+FMA+F16C hosts (Zen 2/3 EPYC, older Xeons).
* operators/avx2/sft_mxfp4-moe.hpp: AVX2_SFT_MXFP4_MOE_TP part on top of
AVX2_MXFP4_MOE_TP (frozen packed E2M1/UE8M0 base, LoRA gate/up/down,
V4 swiglu clamp, transpose-free packed-weight dX, authoritative LoRA grads).
* operators/moe-sft-tp.hpp: AVX2 result merge; AVX-512 merge loops and the
block-FP8 staging guarded so the wrapper builds without AVX-512.
* ext_bindings.cpp: bind MXFP4_SFT_MOE -> TP_MOE_SFT<AVX2_SFT_MXFP4_MOE_TP>
when USE_AMX_AVX_KERNEL is off; MOESFTBindings no longer AMX-only.
* python/sft/backend.py: runtime reports cpu_variant=avx2, kernel avx2-mxfp4-g32.
* test_mxfp4_sft_numerical.py: inference regression uses AVX2MXFP4_MOE when
AMXFP4_KGroup_MOE is absent.
Validated on 2x EPYC 7452 (AVX2, no AVX-512): test_mxfp4_sft_numerical.py
TP1/TP2 all stages rel-L2 <= 0.4% (limits 3-5%); validate_dsv4_mxfp4_sft_layer.py
on DeepSeek-V4-Flash-0731 layer 3, TP1 and TP2: passed, forward/grad_input/
LoRA grads rel-L2 0.10-0.29%, base weights unchanged.
Note: the build uses -ffast-math, so exp() past the float range is undefined;
the scalar sigmoid tails clamp the argument like the vector path.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_011Mgsduzz2fx1gSvfHbR8RZ
* fix(sft): clamp MXFP4 scalar-tail sigmoid exponent
* feat(sft): integrate native V4 experts with LoRA training and export
Wrap only routed experts while preserving V4 routing and shared branches. Validate native checkpoint provenance and stream a BF16 non-expert cache without dequantizing routed experts.
Add audited static SGLang export and reproducible training, resume and generation acceptance helpers. CPU retain/GPU recompute GAS1/GAS4 smoke and 265 KT regression tests pass; long-run quality and resume acceptance remain in progress.
* fix(sft): retain async inference buffers until CUDA work completes
* test(sft): isolate V4 adapter effects with matched-kernel controls
Allow baseline and non-expert-only exports to retain a zero-B expert adapter so all ablations execute the same native forward kernel. Add paired held-out generation reporting with explicitly diagnostic lexical metrics.
* test(sft): compare PEFT targets semantically during V4 resume audits
* test(sft): cancel queued V4 inference checks after a failed request
* test(sft): preserve raw V4 responses before finite-output validation
* feat: prepare V4 cache atomically and package static adapter export
* build: preserve accepted CUDA carriers across Python ABIs
* test: retain FP8 loading contracts after V4 wrapper integration
* test: exclude historical concatenated-sample smoke from delivery
* test(sft): audit unpacked batches and single-GPU resume state
* docs(sft): align April V4 recipe with approved unpacked delivery
* test: keep frozen-parameter audit indices exact for large tensors
* test: record the loaded native backend in acceptance evidence
* test: timestamp training audit events for startup measurements
* docs: pin V4 serving dependencies and account for reserved token slots
* [refactor](sft): keep V4 follow-up focused on runtime integration
---------
Co-authored-by: 05yuki <ceu.azul.samu@gmail.com>
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
* [docs]: add reproducible Kimi K2.5 PyPI LoRA tutorial
* [docs]: highlight Kimi K2.5 and K2.6 LoRA support
* Update README to simplify LoRA fine-tuning entry
* Add updates section to README.md
* [docs]: clarify clean Kimi setup and held-out evaluation
* [docs]: explain pinned model downloads on offline training hosts
* [docs]: isolate Kimi caches in the tutorial work directory
* [docs]: simplify Kimi LoRA installation and training tutorial
* 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>
* 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>
- 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.
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.
* [feat](kt-kernel): add MXFP4 MoE operator with E2M1 weights × BF16 activations
Implements AMX_FP4_MOE_TP based on the RAWINT4 (k2-moe) CRTP pattern.
FP4 E2M1 weights are nibble-packed and decoded via PSHUFB LUT, then
computed with BF16 activations using _mm512_dpbf16_ps. Supports weight-only
per-kgroup scaling (group_size=32) and tensor parallelism.
Includes a Python validation test covering uniform, alternating, ramp,
and random weight patterns.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
* [feat](kt-kernel): adapt MXFP4 MoE backend for DeepSeek-V4-Flash (#1950)
V4-Flash routed experts ship as native MXFP4 (E2M1 nibble + ue8m0 group
scale). Expose AMXFP4_KGroup_MOE through NativeMoEWrapper, add a loader
that handles V4's `layers.{L}.ffn.experts.{i}.{w1,w3,w2}.{weight,scale}`
naming and converts ue8m0 → bf16 via a lossless bit-cast, register the
model entry, and ship an end-to-end numerical validation script.
Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* [perf](kt-kernel): MXFP4 MoE add mat-mat 4×4 tile, refine mat-vec reduce (#1957)
mat_mul_kgroup previously aliased to fp4_mat_vec_kgroup, leaving large
batches stuck on the per-token path. Implement fp4_mat_mat_kgroup as a
4×4 register tile (MB=NB=4, 16 zmm accumulators) so each PSHUFB decode
of four weight rows is reused across four tokens.
Refactor fp4_mat_vec_kgroup to accumulate four N-rows in parallel and
flush them with a new reduce4 helper, removing per-row reduce_add_ps
calls from the hot loop. Mark mxfp4_to_bf16_32 always_inline.
Add bench/bench_fp4_moe.py with --routing {balanced,concentrated} and
a backend registry so future kernels can be added without changing the
runner.
Dispatch thresholds, derived_init, GeneralMOEConfig handling,
load_weights, write_weights_to_buffer and the TP_MOE specialization are
unchanged.
Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* fix(loader): avoid uint16 lshift in ue8m0->bf16 conversion
PyTorch CPU has no lshift kernel for UInt16, so the previous
`(scale_t.to(torch.uint16) << 7)` raised NotImplementedError when
loading any V4-Flash MXFP4 routed-expert scale tensor on the host.
Switch to int32 for the shift (kernel exists) and narrow to int16
afterwards. The shifted value max is 255<<7 = 32640, well within
int16 range, so the narrow is lossless. The .view(bfloat16) bit
pattern is identical (bf16 sign bit is always 0 for ue8m0 values).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* docs(v4-flash): hybrid CPU/GPU recipe + bump kt-sglang submodule
Bumps third_party/sglang to kvcache-ai/sglang main (3cbd49c29) which now
contains DeepSeek V4 Flash model support + consumer-GPU (SM_120) portable
Triton/TileLang fallbacks (kt-sglang PR #38).
Adds doc/en/DeepSeek-V4-Flash.md tutorial: 8x RTX 5090 hybrid recipe with
the full launch command, OpenAI-compatible /generate + /v1/chat/completions
examples, and the kt chat CLI client.
---------
Co-authored-by: ouqingliang <1692110604@qq.com>
Co-authored-by: Claude Opus 4.7 <noreply@anthropic.com>
* [feat]: simplify sglang installation with submodule, auto-sync CI, and version alignment
- Add kvcache-ai/sglang as git submodule at third_party/sglang (branch = main)
- Add top-level install.sh for one-click source installation (sglang + kt-kernel)
- Add sglang-kt as hard dependency in kt-kernel/pyproject.toml
- Add CI workflow to auto-sync sglang submodule daily and create PR
- Add CI workflow to build and publish sglang-kt to PyPI
- Integrate sglang-kt build into release-pypi.yml (version.py bump publishes both packages)
- Align sglang-kt version with ktransformers via SGLANG_KT_VERSION env var injection
- Update Dockerfile to use submodule and inject aligned version
- Update all 13 doc files, CLI hints, and i18n strings to reference new install methods
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* [build]: bump version to 0.5.2
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* [build]: rename PyPI package from kt-kernel to ktransformers
Users can now `pip install ktransformers` to get everything
(sglang-kt is auto-installed as a dependency).
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* Revert "[build]: rename PyPI package from kt-kernel to ktransformers"
This reverts commit e0cbbf6364.
* [build]: add ktransformers meta-package for PyPI
`pip install ktransformers` now works as a single install command.
It pulls kt-kernel (which in turn pulls sglang-kt).
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* [fix]: show sglang-kt package version in kt version command
- Prioritize sglang-kt package version (aligned with ktransformers)
over sglang internal __version__
- Update display name from "sglang" to "sglang-kt"
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* [fix]: improve sglang-kt detection in kt doctor and kt version
Recognize sglang-kt package name as proof of kvcache-ai fork installation.
Previously both commands fell through to "PyPI (not recommended)" for
non-editable local source installs. Now version.py reuses the centralized
check_sglang_installation() logic.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* [build]: bump version to 0.5.2.post1
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
yyj