* [feat](kt-kernel): add AVX2 RAWINT4 routed-expert LoRA SFT backend
* [fix](kt-kernel): fix RAWINT4 SFT wrapper and clamp gradients
Allow the RAWINT4 production wrapper on AVX2 runtimes and reject shared backward pools before construction. Preserve FP32 clamp decisions across BF16 caches, remove unused configuration state, and cover rectangular dimensions, multiblock dX, clamp boundaries, and production wrapper tasks with tighter LoRA gradient checks.
* [fix](kt-kernel): plumb swiglu_limit through the RAWINT4 SFT python path
The production wrapper hard-zeroed swiglu_limit for every format except MXFP4, and both SFT python gates rejected non-zero values for RAWINT4 even though MOESFTConfig carries the field and the AVX2 kernel implements the clamp in forward plus the gradient mask in backward. RAWINT4_SFT now accepts a finite non-negative swiglu_limit end to end, while the AMX K2 part (which computes plain SiLU) fails fast instead of silently dropping the clamp. The production-wrapper contract test is parametrized over swiglu_limit 0 and 2, a model-config plumb test covers plain and nested text_config shapes, the __main__ entry exercises both limits, and bench_rawint4_sft.py gains --swiglu-limit.
* 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>
* 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>
* [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.
* [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.
* [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>
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>
* [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>
Add numa_nodes parameter to BaseMoEWrapper and all subclasses, allowing
users to explicitly specify which NUMA node IDs to use for subpool
mapping instead of always defaulting to sequential [0, 1, ..., N-1].
This enables running multiple KTransformers instances on different NUMA
nodes of the same machine, e.g. --kt-threadpool-count 1 --kt-numa-nodes 1
to bind to NUMA node 1. Previously this required external numactl
workarounds since subpool_numa_map was hardcoded to start from 0.
* Increase timeout for Check if --kt-gpu-prefill-token-threshold is in the help output to 90 seconds.
In cloud environments,CUDA initialization and Python module loading can easily exceed 30 seconds.
* Update kt-kernel/python/cli/utils/sglang_checker.py
add comment about the change
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>
* [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>
* Handle unquoted paths and special characters in model scanner
* Fix ValueError: capture_output cannot be used with stderr
`capture_output=True` internally sets `stderr=PIPE`, which conflicts
with `stderr=subprocess.DEVNULL`. Replace `capture_output=True` with
explicit `stdout=subprocess.PIPE` to keep stderr suppressed correctly.
Also remove redundant `shell=False` (already the default).
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---------
Co-authored-by: ErvinXie <ervinxie@foxmail.com>
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* [feat]: redesign kt run interactive configuration with i18n support
- Redesign kt run with 8-step interactive flow (model selection, inference method, NUMA/CPU, GPU experts, KV cache, GPU/TP selection, parsers, host/port)
- Add configuration save/load system (~/.ktransformers/run_configs.yaml)
- Add i18n support for kt chat (en/zh translations)
- Add universal input validators with auto-retry and Chinese comma support
- Add port availability checker with auto-suggestion
- Add parser configuration (--tool-call-parser, --reasoning-parser)
- Remove tuna command and clean up redundant files
- Fix: variable reference bug in run.py, filter to show only MoE models
* [feat]: unify model selection UI and enable shared experts fusion by default
- Unify kt run model selection table with kt model list display
* Add Total size, MoE Size, Repo, and SHA256 status columns
* Use consistent formatting and styling
* Improve user decision-making with more information
- Enable --disable-shared-experts-fusion by default
* Change default value from False to True
* Users can still override with --enable-shared-experts-fusion
* [feat]: improve kt chat with performance metrics and better CJK support
- Add performance metrics display after each response
* Total time, TTFT (Time To First Token), TPOT (Time Per Output Token)
* Accurate input/output token counts using model tokenizer
* Fallback to estimation if tokenizer unavailable
* Metrics shown in dim style (not prominent)
- Fix Chinese character input issues
* Replace Prompt.ask() with console.input() for better CJK support
* Fixes backspace deletion showing half-characters
- Suppress NumPy subnormal warnings
* Filter "The value of the smallest subnormal" warnings
* Cleaner CLI output on certain hardware environments
* [fix]: correct TTFT measurement in kt chat
- Move start_time initialization before API call
- Previously start_time was set when receiving first chunk, causing TTFT ≈ 0ms
- Now correctly measures time from request sent to first token received
* [docs]: 添加 Clawdbot 集成指南 - KTransformers 企业级 AI 助手部署方案
* [docs]: 强调推荐使用 Kimi K2.5 作为核心模型，突出企业级推理能力
* [docs]: 添加 Clawdbot 飞书接入教程链接
* [feat]: improve CLI table display, model verification, and chat experience
- Add sequence number (#) column to all model tables by default
- Filter kt edit to show only MoE GPU models (exclude AMX)
- Extend kt model verify to check *.json and *.py files in addition to weights
- Fix re-verification bug where repaired files caused false failures
- Suppress tokenizer debug output in kt chat token counting
* [fix]: fix cpu cores.
---------
Co-authored-by: skqliao <skqliao@gmail.com>
losewayy