feat: DeepSeek-V4-Flash integration + Blackwell (sm120) bring-up - #153

Draft

drunkcoding wants to merge 96 commits into

mainfrom

feature/deepseek-v4-kernel-integration

Commits on May 12, 2026

add gitignore
leyang.xue
committed
02bba41

Commits on May 21, 2026

test: add shared kernel test infrastructure (conftest + package init)

leyang.xue
and
sisyphus-dev-ai
committed
dc46551
feat: add per-head fused QK RMSNorm Triton kernel for V4 MLA

leyang.xue
and
sisyphus-dev-ai
committed
a395435
feat: add native CUDA fused QK RMSNorm for V4 MLA (Path A)

leyang.xue
and
sisyphus-dev-ai
committed
7033852

Commits on May 24, 2026

feat(kernels): consolidate v4 kernels into batchgen_kernels

leyang.xue
and
sisyphus-dev-ai
committed
3cf542a
feat(kernels): v4 DSA attention kernel improvements

leyang.xue
and
sisyphus-dev-ai
committed
e0bb7dc
test: add v4 kernel test suite

leyang.xue
and
sisyphus-dev-ai
committed
5d9e63e
feat(model): V4 Flash model scaffolding + HC kernel wiring

leyang.xue
and
sisyphus-dev-ai
committed
049b0a3
chore: docker base image update + vllm entry points doc

leyang.xue
and
sisyphus-dev-ai
committed
b31d003
fix(kv-cache): correct V4 KV profile to 584 bytes/token

leyang.xue
and
sisyphus-dev-ai
committed
6994ad6
feat(model): wire V4 gate routing, MoE activation, and attention decode

leyang.xue
and
sisyphus-dev-ai
committed
25998e2
test: add V4 wiring correctness tests and fix HC references

leyang.xue
and
sisyphus-dev-ai
committed
51c3127

Commits on Jun 1, 2026

fix(v4flash): make DP-decode MoE collective symmetric and correct
leyang.xue
committed
3afd3ae

Commits on Jun 2, 2026

fix(v4flash): wire c4 indexer DP-gather + functional compressor for decode
leyang.xue
committed
9628082
fix(v4flash): pass compressed seqlen to c4 sparse indexer (fix decode IMA)
leyang.xue
committed
d3d129a
fix(v4flash): gather decode results from global_batch, not local maps
leyang.xue
committed
cdc5c1a

Commits on Jun 4, 2026

feat(v4flash): add sm120 Triton kernels for grouped MoE and sparse MLA decode

leyang.xue
and
sisyphus-dev-ai
committed
92b4fb5
feat(v4flash): add DeepSeek-V4 single KV pool and decode KV coordinator

leyang.xue
and
sisyphus-dev-ai
committed
13574ff
feat(v4flash): add sparse MLA decode adapter, torch reference, and prefill populate

leyang.xue
and
sisyphus-dev-ai
committed
5bd3250
feat(v4flash): wire sm120 decode path with grouped MoE, persistent experts, and PyNCCL collectives

leyang.xue
and
sisyphus-dev-ai
committed
e86e584
feat(v4flash): add checkpoint metadata loader and assets packaging

leyang.xue
and
sisyphus-dev-ai
committed
f462863
fix(v4flash): register tokenizer discovery for V4-Flash

leyang.xue
and
sisyphus-dev-ai
committed
7945f3c
chore(models): align v3/glm5/minimax/kimi config with shared decode changes

leyang.xue
and
sisyphus-dev-ai
committed
1a62675
chore(kv-cache): update host KV manager config

leyang.xue
and
sisyphus-dev-ai
committed
6b0e660
feat(server): wire decode-timing step_done and worker/process management

leyang.xue
and
sisyphus-dev-ai
committed
bea1c26
feat(core): expose get_tensor and weight-buffer support for persistent experts

leyang.xue
and
sisyphus-dev-ai
committed
4be62c2
build(kernels): update sm120 kernel build, op builder, and docker image

leyang.xue
and
sisyphus-dev-ai
committed
9bbb050
test(v4flash): add V4 decode integration, KV coordinator, and compressor tests

leyang.xue
and
sisyphus-dev-ai
committed
d7a1874
test(e2e): add serving benchmark and cross-framework MMLU-Pro harness

leyang.xue
and
sisyphus-dev-ai
committed
4d765cf
chore(tools): add V4 repro, divtrace, and MoE analysis scripts

leyang.xue
and
sisyphus-dev-ai
committed
6628d51

Commits on Jun 6, 2026

chore: ignore generated batch storage artifacts

leyang.xue
and
sisyphus-dev-ai
committed
9ce9af0
feat(v4flash): support prepacked multi-seq prefill slicing

leyang.xue
and
sisyphus-dev-ai
committed
41ca4bf
style(server): normalize worker main loop indentation

leyang.xue
and
sisyphus-dev-ai
committed
de030c4
feat(kernels): retarget sm120 JIT extensions for Blackwell

leyang.xue
and
sisyphus-dev-ai
committed
a081227
feat(kv-cache): isolate host KV shared memory names

leyang.xue
and
sisyphus-dev-ai
committed
0d6da90
feat(server): expose prepack and memory tuning flags

leyang.xue
and
sisyphus-dev-ai
committed
faa39e2
fix(parameter-server): support V4 remote weight serving

leyang.xue
and
sisyphus-dev-ai
committed
062b947
feat(v4flash): add grouped MXFP4 MoE decode path

leyang.xue
and
sisyphus-dev-ai
committed
a67fd17
perf(v4flash-attn): vectorize sparse index physicalization

leyang.xue
and
sisyphus-dev-ai
committed
2ccdfda
test(e2e): add V4 MMLU-Pro batch harness

leyang.xue
and
sisyphus-dev-ai
committed
3b4e323
perf(v4flash-attn): batch C128 decode metadata updates

leyang.xue
and
sisyphus-dev-ai
committed
acb143c

Commits on Jun 9, 2026

Commits on Jun 15, 2026

feat(ckpt-converter): add load_rank_shard_tensors for selective rank-shard loading

leyang.xue
and
sisyphus-dev-ai
committed
7f69ce1
fix(v4flash): correct fp8 cast in act_quant and add shared-expert swiglu_limit

leyang.xue
and
sisyphus-dev-ai
committed
3aef181
feat(v4flash): add sparse prefill attention ported from official reference

leyang.xue
and
sisyphus-dev-ai
committed
4ae7285
feat(v4flash): wire sparse prefill routing, opt-in QAT linear, and fix MoE collective backend

leyang.xue
and
sisyphus-dev-ai
committed
2fd7f3d
fix(v4flash): use compression-aware rope cache for sparse and compressed layers

leyang.xue
and
sisyphus-dev-ai
committed
7405257
feat(v4flash-attn): add tensor-dump diagnostics and guard page-table fast path

leyang.xue
and
sisyphus-dev-ai
committed
a987091
fix(server): make NCCL timeout configurable for long-sequence offload decode

leyang.xue
and
sisyphus-dev-ai
committed
fcd289a
fix(v4flash): keep prefill experts streamed and load rank-local vocab-sharded roots

leyang.xue
and
sisyphus-dev-ai
committed
722898d
fix(kv-cache): charge V4 KV pools in compressed token space to prevent page exhaustion

leyang.xue
and
sisyphus-dev-ai
committed
9d46f71
fix(v4flash): join vocab-parallel collectives on 0-seq ranks and use pool-aware admission

leyang.xue
and
sisyphus-dev-ai
committed
3d408fa
chore(tools): add V4 MMLU eval, divtrace, and sanity scripts

leyang.xue
and
sisyphus-dev-ai
committed
944a57e
docs: add DeepSeek-V4-Flash Blackwell bring-up handoff notes

leyang.xue
and
sisyphus-dev-ai
committed
2ab2c53
fix(v4flash): make ON_HOLD eviction V4-coordinator-aware to avoid crash

leyang.xue
and
sisyphus-dev-ai
committed
12e2e2a
fix(v4flash): collectively re-init GPU KV coordinator on watermark re-prefill

leyang.xue
and
sisyphus-dev-ai
committed
3db3385
docs: record watermark on-hold/re-prefill crash fixes and QAT char-exact plan

leyang.xue
and
sisyphus-dev-ai
committed
8845d37
docs: localize residual decode drift to non-QAT grouped MoE kernel

leyang.xue
and
sisyphus-dev-ai
committed
bce0b40
test(v4flash): add grouped MoE kernel vs per-expert parity test (xfail)

leyang.xue
and
sisyphus-dev-ai
committed
986331e
feat(v4flash): add QAT-faithful grouped MoE decode kernel (BATCHGEN_V4_QAT_MOE)

leyang.xue
and
sisyphus-dev-ai
committed
fdca802
docs: record QAT MoE kernel result and residual #2 (decode attention/router)

leyang.xue
and
sisyphus-dev-ai
committed
22c8686
docs: narrow residual #2 (rule out SM120 kernel + rope mismatch)

leyang.xue
and
sisyphus-dev-ai
committed
e72f1d5
docs: refute fp8 KV, localize residual #2 to sparse indexer scoring

leyang.xue
and
sisyphus-dev-ai
committed
ee077d2
docs: validate indexer missing-relu changes ~30% of topk KV selection

leyang.xue
and
sisyphus-dev-ai
committed
5d61f23
fix(v4flash): apply per-head relu in sparse indexer score to match official

leyang.xue
and
sisyphus-dev-ai
committed
42fd462
fix(v4flash): fp4-fake-quant indexer q and k to match official QAT

leyang.xue
and
sisyphus-dev-ai
committed
cba7f54
docs: record indexer relu (major) + fp4 (marginal) char-exact results

leyang.xue
and
sisyphus-dev-ai
committed
388da19
docs: root-cause residual to inherent fp rounding (router + sink refuted)

leyang.xue
and
sisyphus-dev-ai
committed
30ffb98

Commits on Jun 16, 2026

build(docker): add GPU_ARCH=hopper support to batchgen image
leyang.xue
committed
bacba50

Commits on Jun 23, 2026

chore(v4flash): drop stale handoff/diagnostic docs and tooling
leyang.xue
committed
0143b23
perf(v4flash): vectorize sm120 decode index builders
leyang.xue
committed
a579e7f
perf(v4flash): warmup, pin autotune, and guard cache geometry for sm120 sparse decode
leyang.xue
committed
813c791
perf(v4flash): vectorize fp8 KV-cache pack (7-tile loop -> batched)
leyang.xue
committed
c911974
test(v4flash): fix stale wkv refs and add eager<->kernel compressor parity
leyang.xue
committed
7f10e5e
feat(v4flash): sm-aware indexer quant (MXFP4 on sm120, FP8 on sm90)
leyang.xue
committed
4a594b3
feat(v4flash): gate grouped MXFP4 MoE off below sm120
leyang.xue
committed
9c558bf
test(v4flash): skip-guard MXFP4-only tests below sm120 + add fp8/dispatch coverage
leyang.xue
committed
cdfc8a8
test(v4flash): fix stale wkv/world_size refs + tilelang guards in decode-loop e2e
leyang.xue
committed
95a0c3e
docs(docker): add H20 rebuild+launch runbook for full V4-Flash run
leyang.xue
committed
e19e464

Commits on Jun 24, 2026

build(docker): add CN-mirror build args + harden GitHub clones
leyang.xue
committed
ba22d80
build(docker): wire CN-mirror build args into H20 runbook script
leyang.xue
committed
cc90d02

Commits on Jun 25, 2026

fix(v4flash): make DP-decode entry collectives rank-safe
leyang.xue
committed
e4d84dc
feat(v4flash): add BATCHGEN_DECODE_DEADLOCK_TRACE rank markers
leyang.xue
committed
8dd373a
perf(v4flash): batch per-expert host sync in _run_owned_experts
leyang.xue
committed
e6a4988
build(docker): persist torch JIT cache + fix memlock for V4 H20 run
leyang.xue
committed
5000c7e
build(deps): bump 7 deps to patch Dependabot vulnerabilities
leyang.xue
committed
ea79238

Commits on Jul 4, 2026

feat(v4flash): add sm120 grouped-MoE kernels (mega3/mega/graph/ragged) + FP4 utils
leyang.xue
committed
d02be92
feat(v4flash): wire mega3 MoE + resident-experts/EP-offload serving path
leyang.xue
committed
9bf47d8
fix(v4flash): size RoPE caches to original_seq_len, grow on demand
leyang.xue
committed
7c25580
fix(v4flash): GPU-KV admission backpressure for prefill and decode
leyang.xue
committed
c8cce99
test(v4flash): numerics parity, arch gate, and MoE wiring parity tests
leyang.xue
committed
14d4593
build(docker): update H20 runbook for sm120 kernel integration
leyang.xue
committed
ab125aa
docs(v4flash): 4xRTX6000-Pro setup guide, kernel/EP notes, autoresearch harness
leyang.xue
committed
f824046
feat(planner): env overrides for phase/module micro-batching knobs
leyang.xue
committed
acf8469