Skip to content
Draft
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
181 commits
Select commit Hold shift + click to select a range
e194e03
docs(deepseek-v4): add development plan and architecture deep dive
wenxie-amd Apr 28, 2026
d3383c0
feat(deepseek-v4)[P1]: add model + training yaml configs
wenxie-amd Apr 28, 2026
8ae1000
feat(deepseek-v4)[P2]: register model_type=deepseek_v4 dispatch path
wenxie-amd Apr 28, 2026
a5d2a56
feat(deepseek-v4)[P3]: layer specs + transformer block scaffolding
wenxie-amd Apr 28, 2026
3b7ad8c
feat(deepseek-v4)[P4]: HC + Hybrid Attention (Dense/HCA/CSA + sink + …
wenxie-amd Apr 28, 2026
5e4008d
feat(deepseek-v4)[P5]: V4 MoE (hash + sqrtsoftplus) + clamped SwiGLU …
wenxie-amd Apr 28, 2026
97b9720
feat(deepseek-v4)[P6-P7]: complete distributed bring-up on single node
wenxie-amd Apr 29, 2026
1030293
docs(deepseek-v4): add phase 8+ replan docs and reference notes
wenxie-amd Apr 29, 2026
df273a4
feat(deepseek-v4)[P8]: root runtime path on LanguageModule spec tree
wenxie-amd Apr 29, 2026
e5fec96
feat(deepseek-v4)[P9]: providerize TE/Turbo module paths
wenxie-amd Apr 30, 2026
b38e83c
feat(deepseek-v4)[P10]: enforce MoE provider path and add V4 config s…
wenxie-amd Apr 30, 2026
752b753
fix(deepseek-v4)[P10]: stabilize smoke runtime and add phase report
wenxie-amd Apr 30, 2026
636ab3d
docs(deepseek-v4)[P12]: plan-2 lockdown + as-built techblog + roadmap…
wenxie-amd May 1, 2026
cad0fb3
feat(deepseek-v4)[P13]: rebase V4 attention on MLASelfAttention (fait…
wenxie-amd May 1, 2026
aa9929a
feat(deepseek-v4)[P13]: fold Compressor/Indexer + TP-shard projection…
wenxie-amd May 1, 2026
1a8bf32
feat(deepseek-v4)[P14]: faithful pre-mul clamped SwiGLU + V4 routers …
wenxie-amd May 1, 2026
5fe8bc3
feat(deepseek-v4)[P14]: V4 MoE structural bring-up + G5 numerical ali…
wenxie-amd May 1, 2026
25ccdb5
feat(deepseek-v4)[P15]: V4 hybrid layer / block subclass refactor + t…
wenxie-amd May 1, 2026
6c5875d
feat(deepseek-v4)[P16]: spec-based MTP integration via upstream Multi…
wenxie-amd May 1, 2026
f548d8b
docs(deepseek-v4 plan-2): defer HF checkpoint adapter to P22+, repurp…
wenxie-amd May 1, 2026
e591b89
refactor(deepseek-v4 P17): retire legacy MTP block + dedup _RMSNorm +…
wenxie-amd May 1, 2026
b583267
refactor(deepseek-v4 P18): provider singleton + activation_func consi…
wenxie-amd May 1, 2026
e66492d
docs(deepseek-v4): replace status.md commit placeholders with concret…
wenxie-amd May 1, 2026
83c33ad
feat(deepseek-v4)[P19]: PP wire-shape + hash-router input_ids upfront…
wenxie-amd May 7, 2026
dba2716
docs(deepseek-v4 plan-2): close out plan-2 status — P19 wrap-up + pla…
wenxie-amd May 7, 2026
79071ad
docs(deepseek-v4 plan-3): scope reporting fixes + Primus-Turbo enable…
wenxie-amd May 7, 2026
4c27787
feat(deepseek-v4)[P20]: V4-aware num_floating_point_operations (close…
wenxie-amd May 7, 2026
ad29221
docs(deepseek-v4)[P20]: pin status.md commit cells to the P20 SHA (4c…
wenxie-amd May 7, 2026
a4419ac
feat(deepseek-v4)[P21]: strict V4 spec build (no nn.Linear fallback)
wenxie-amd May 7, 2026
5573703
docs(deepseek-v4)[P21]: pin status.md / plan-3 commit cells to the P2…
wenxie-amd May 7, 2026
4c362bc
feat(deepseek-v4)[P22]: core_attention submodule for V4 dense layers …
wenxie-amd May 8, 2026
8468459
docs(deepseek-v4)[P22]: pin status.md / plan-3 commit cells to the P2…
wenxie-amd May 8, 2026
6d18500
feat(deepseek-v4)[P23]: Turbo DeepEP dispatcher in V4 layer specs
wenxie-amd May 8, 2026
04e1eb3
docs(deepseek-v4)[P23]: pin status.md / plan-3 commit cells to the P2…
wenxie-amd May 8, 2026
f467a73
docs(deepseek-v4)[plan-4]: open plan-4 — in-tree Triton-fused V4 atte…
wenxie-amd May 8, 2026
8b97188
feat(deepseek-v4)[P24]: extract V4 attention math into v4_attention_k…
wenxie-amd May 8, 2026
2b12d67
docs(deepseek-v4)[P24]: pin status.md commit cells to the P24 SHA (8b…
wenxie-amd May 8, 2026
38ef526
refactor(deepseek-v4)[P24]: keep V4 eager attention matmuls on bf16 t…
wenxie-amd May 8, 2026
0c19780
docs(deepseek-v4)[P24]: pin status.md commit cell to the P24 dtype-co…
wenxie-amd May 8, 2026
7df4cfe
feat(deepseek-v4)[P25]: in-tree Primus Triton v4_attention kernel for…
wenxie-amd May 8, 2026
1ba38ba
docs(deepseek-v4)[P25]: pin status.md commit cells to the P25 SHA (7d…
wenxie-amd May 8, 2026
da6f48b
feat(deepseek-v4)[P26]: in-tree Primus Triton v4_csa_attention kernel…
wenxie-amd May 8, 2026
36dfca6
docs(deepseek-v4)[P26]: pin status.md commit cells to the P26 SHA (da…
wenxie-amd May 8, 2026
6a17c3b
docs(deepseek-v4)[plan-4]: reorder P27 — release-tier shape gate (G28…
wenxie-amd May 8, 2026
e19663f
feat(deepseek-v4)[plan-4][P27]: release-tier shape gate + dispatch
wenxie-amd May 9, 2026
03bacc2
docs(deepseek-v4)[P27]: pin status.md + plan-4 hand-off cells to the …
wenxie-amd May 9, 2026
578496b
docs(deepseek-v4)[plan-5]: open plan-5 (perf optimisation, P28..P32) …
wenxie-amd May 9, 2026
afd7ea5
feat(deepseek-v4)[plan-5][P28]: V4-Flash proxy + EP=8 baseline trace …
wenxie-amd May 9, 2026
2ec7e40
docs(deepseek-v4)[P28]: pin status.md P28 cells to the P28 SHA (afd7e…
wenxie-amd May 9, 2026
1ea7e7a
feat(deepseek-v4)[plan-5][P29]: torch.compile-fused HyperMixer Sinkho…
wenxie-amd May 9, 2026
e653aac
docs(deepseek-v4)[P29]: pin status.md P29 cells to the P29 SHA (1ea7e…
wenxie-amd May 9, 2026
e8ca33e
feat(deepseek-v4)[plan-5][P30]: prune V4 dense and HCA attention SWA …
wenxie-amd May 9, 2026
d4d4aee
docs(deepseek-v4)[P30]: pin status.md P30 cells to the P30 SHA (e8ca3…
wenxie-amd May 9, 2026
2c7cf59
feat(deepseek-v4)[plan-5][P31]: split CSA backward into dense local a…
wenxie-amd May 9, 2026
945ec74
docs(deepseek-v4)[P31b]: add split CSA backward profile report
wenxie-amd May 9, 2026
a3f5f46
docs(deepseek-v4)[plan-5]: remove Phase 32 planning
wenxie-amd May 9, 2026
e5987f9
feat(deepseek-v4)[plan-5][P32]: split CSA FWD + atomic-free V4/CSA BW…
wenxie-amd May 11, 2026
ded5348
feat(deepseek-v4)[plan-5][P32]: ship CSA FWD split; gate proxy-regres…
wenxie-amd May 11, 2026
284b805
docs(deepseek-v4)[plan-5][P32]: record A/B verification confirming pr…
wenxie-amd May 14, 2026
b03416d
feat(deepseek-v4)[plan-5][P32]: dual-RoPE bf16 cast fix + flip split/…
wenxie-amd May 14, 2026
ad81766
docs(deepseek-v4)[plan-6]: open plan-6 (elemwise / layout fusion + TF…
wenxie-amd May 14, 2026
79f10df
feat(deepseek-v4)[plan-6][P33]: V4 TFLOP/s closed-form correction (SW…
wenxie-amd May 14, 2026
ac8a54e
docs(deepseek-v4)[plan-6][P33]: pin phase 33 commit SHAs in status.md
wenxie-amd May 14, 2026
29baf15
feat(deepseek-v4)[plan-6][P34]: _stack_grouped_linear_weight Triton F…
wenxie-amd May 14, 2026
a9c975c
docs(deepseek-v4)[plan-6][P34]: pin phase 34 commit SHAs in status.md
wenxie-amd May 14, 2026
842110f
plan-6 P35: fuse apply_interleaved_partial_rope FWD/BWD into Triton k…
wenxie-amd May 14, 2026
2ae9914
plan-6 P35: pin status.md to the P35 commit (842110f8)
wenxie-amd May 14, 2026
ed44bc8
plan-6 P36: fuse sinkhorn_normalize FWD/BWD into Triton kernel
wenxie-amd May 14, 2026
ff930b4
plan-6 P36: pin status.md to the P36 commit (ed44bc87)
wenxie-amd May 14, 2026
f451223
backend(deepseek-v4)[plan-6][P37]: HyperConnection compute_weights ta…
wenxie-amd May 14, 2026
879ca6f
backend(deepseek-v4)[plan-6][P38]: Indexer.forward scoring Triton FWD…
wenxie-amd May 14, 2026
7f39e02
backend(deepseek-v4)[plan-6][P39]: V4 router post-logits Triton FWD/B…
wenxie-amd May 14, 2026
041170a
plan-6 P39: pin status.md and p39-summary.md to the P39 commit (7f39e…
wenxie-amd May 14, 2026
b08975b
docs(deepseek-v4)[plan-6][P40]: plan-6 close-out -- proxy bake-off + …
wenxie-amd May 14, 2026
6a5c2b5
plan-6 P40: pin status.md and p40-summary.md to the P40 close-out com…
wenxie-amd May 14, 2026
2c73e73
docs(deepseek-v4)[plan-6][P40]: P40 close-out trace + R2.6 rule + 202…
wenxie-amd May 14, 2026
c7eaad0
docs(deepseek-v4)[plan-6][plan-7]: extend plan-6 with P41/P42/P43/P44…
wenxie-amd May 15, 2026
9444f8a
feat(deepseek-v4)[plan-6][P41]: Indexer.forward post-einsum tail Trit…
wenxie-amd May 15, 2026
78ddbf7
plan-6 P41: pin status.md and p41-summary.md to the P41 commit (9444f…
wenxie-amd May 15, 2026
9898aae
docs(deepseek-v4)[plan-6][P42]: descope at task-list refinement (R9.1…
wenxie-amd May 15, 2026
68be37c
plan-6 P42: pin status.md and p42-summary.md to the P42 descope commi…
wenxie-amd May 15, 2026
77ab0ca
docs(deepseek-v4)[plan-6][P43]: router post-logits 50-iter 3-run A/B …
wenxie-amd May 15, 2026
1244f3f
plan-6 P43: pin status.md and p43-summary.md to the P43 commit (77ab0…
wenxie-amd May 15, 2026
15b680c
docs(deepseek-v4)[plan-6][P44]: descope at task-list refinement -- at…
wenxie-amd May 15, 2026
fd26f70
plan-6 P44: pin status.md and p44-summary.md to the P44 descope commi…
wenxie-amd May 15, 2026
d925c56
feat(deepseek-v4)[plan-7][P45]: multi-tensor BF16 add Triton prototyp…
wenxie-amd May 15, 2026
c7fe33b
plan-7 P45: pin status.md and p45-summary.md to the P45 commit (d925c…
wenxie-amd May 15, 2026
0f95e81
docs(deepseek-v4)[plan-7][P46-P48]: P46/P47 descope via P45 evidence;…
wenxie-amd May 15, 2026
a1cc391
plan-7 P46/P47/P48: pin status.md + p46/p47/p48-summary.md to the pla…
wenxie-amd May 15, 2026
f4d63e5
docs(deepseek-v4)[plan-8]: open plan-8 with tilelang V4 attention sco…
wenxie-amd May 15, 2026
73f763e
feat(deepseek-v4)[plan-8][P49]: tilelang infra + dispatcher (default-…
wenxie-amd May 15, 2026
9f885bb
plan-8 P49: pin status.md and p49-summary.md to the P49 commit (73f76…
wenxie-amd May 15, 2026
af07de9
feat(deepseek-v4)[plan-8][P50]: V4 dense FWD tilelang kernel cr=0 (de…
wenxie-amd May 15, 2026
e13b21a
plan-8 P50: pin status.md and p50-summary.md to the P50 commit (af07d…
wenxie-amd May 15, 2026
ed654ca
feat(deepseek-v4)[plan-8][P51]: V4 dense BWD tilelang + autograd wrap…
wenxie-amd May 15, 2026
9ae6df0
plan-8 P51: pin status.md and p51-summary.md to the P51 commit (ed654…
wenxie-amd May 15, 2026
19e41c9
docs(deepseek-v4)[plan-8][P52-P56]: HCA/CSA tilelang descope (V4-Flas…
wenxie-amd May 15, 2026
28a6766
plan-8 P52-P56: pin status.md + p52..p56-summary.md to the plan-8 clo…
wenxie-amd May 15, 2026
112d0dd
docs(deepseek-v4)[plan-8][P57]: open P57 Triton V4 attention perf pus…
wenxie-amd May 15, 2026
1e4bb2e
feat(deepseek-v4)[plan-8][P57]: Round-1 integration cr=0 3.01 / cr=4 …
wenxie-amd May 15, 2026
fa67a9f
feat(deepseek-v4)[plan-8][P57] Round-1 integration: cr=0 BWD 3.01 / c…
wenxie-amd May 15, 2026
048e225
feat(deepseek-v4)[plan-8][P57] R2 close-out: cr=0/4/128 V4 attention …
wenxie-amd May 15, 2026
7ce3444
docs(deepseek-v4)[plan-8][P57]: status-pin Phase 57 R2 close-out rows…
wenxie-amd May 15, 2026
9bd9cd1
feat(deepseek-v4)[plan-8][P57]: close-out 2 — tilelang optional + CLI…
wenxie-amd May 15, 2026
86b074d
docs(deepseek-v4)[plan-8][P57]: status-pin Phase 57 close-out 2 row t…
wenxie-amd May 15, 2026
5caf08e
feat(deepseek-v4): add Pro + Muon single-node launcher (#729)
wenxie-amd May 21, 2026
f3587ca
fix(deepseek-v4): correct NUM_STAGES_FWD to solve the NaN issue
lhzhang333 Jun 8, 2026
d897076
fix(deepseek-v4): add the batch axis in sin/cos in compressed pool to…
lhzhang333 Jun 8, 2026
6d353b4
feat(deepseek-v4): recompte compatibility
lhzhang333 Jun 8, 2026
38f3bdc
update script
lhzhang333 Jun 9, 2026
aa7aa65
freeze indexer param to enable overlap_grad_reduce/param_gather
lhzhang333 Jun 11, 2026
3841e36
update script: enable overlap_grad_reduce/param_gather and manual_gc
lhzhang333 Jun 11, 2026
710d94d
feat(deepseek-v4): enable faithful end-to-end MTP with per-depth Hype…
lhzhang333 Jun 11, 2026
3c25e44
feat(deepseek-v4): add FP8 (E4M3) Indexer QK path for the CSA selector
lhzhang333 Jun 11, 2026
e9d17fd
feat(deepseek-v4): wire Muon with DeepSeek-V4 hybrid Newton-Schulz + …
lhzhang333 Jun 11, 2026
d463c7b
feat(deepseek-v4): launch plumbing for Muon/dist_muon + MXFP8 (+ fp8 …
lhzhang333 Jun 11, 2026
e7c70a7
update script: add pp-warmup
lhzhang333 Jun 12, 2026
689a260
update script: remove extra comment
lhzhang333 Jun 12, 2026
756b316
add dsv4 flash script
lhzhang333 Jun 17, 2026
80ac892
enable mtp layer and fp8 indexer
lhzhang333 Jun 17, 2026
f45a8a3
feat(deepseek-v4): auto-install emerging_optimizers + quiet Muon/Trit…
wenxie-amd Jun 18, 2026
dac0a60
update script: add manual_gc to pro muon run
wenxie-amd Jun 18, 2026
f9b924f
feat(deepseek-v4): add training performance projection toolkit + site
wenxie-amd Jun 18, 2026
128f934
fix(deepseek-v4 projection): capture all cr with dist-opt off; track …
wenxie-amd Jun 18, 2026
1384067
feat(deepseek-v4 projection): real flash data + fix attn fwd/bwd phase
wenxie-amd Jun 18, 2026
f5b616e
feat(deepseek-v4 projection): calibrate iter time + analytic FLOPs to…
wenxie-amd Jun 18, 2026
824f0b7
feat(deepseek-v4 projection site): user-set world size + default/layo…
wenxie-amd Jun 19, 2026
76bb62a
opt(deepseek-v4 projection): fix trace attribution, loss bucketing, c…
lhzhang333 Jun 22, 2026
4a90e16
fix(deepseek-v4 projection): correct capture meta + param count; clar…
wenxie-amd Jun 22, 2026
4b8b9a3
opt(deepseek-v4 projection): re-capture flash traces cleanly + recali…
lhzhang333 Jun 22, 2026
67b6b7e
ci(deepseek-v4 projection): deploy from dev/tas/deepseek-v4 branch
wenxie-amd Jun 22, 2026
6155507
DSv4: support in the performance-projection benchmark (#776)
JohnQinAMD Jun 23, 2026
0cfd6ba
DSv4: FP8/FP4 low-precision training recipe + Muon mxfp8 fix + launch…
JohnQinAMD Jun 23, 2026
e89b4c0
DSv4: FlyDSL attention backend + MI355X Triton tuning (#778)
JohnQinAMD Jun 23, 2026
aaa4f27
Dev/john/dsv4 gfx1250 port (#785)
JohnQinAMD Jun 25, 2026
159c196
feat: bump primus turbo wrapper to latest (#786)
RuibinCheung Jun 25, 2026
cb9510d
v4-pro launcher: default-ON the tested Triton small-kernel fusions
JohnQinAMD Jun 25, 2026
135ec6f
feat: add pad_offsets to match latest TE
RuibinCheung Jun 25, 2026
574b9aa
fix(deepseek-v4 projection): correct optimizer sharding model
lhzhang333 Jun 25, 2026
5b09425
opt(deepseek-v4 projection): model MTP flops and timing
lhzhang333 Jun 25, 2026
2390d56
opt(deepseek-v4 projection): support pipeline layout and selective re…
lhzhang333 Jun 25, 2026
cdba154
fix(deepseek-v4 projection): improve trace attribution and capture pa…
lhzhang333 Jun 25, 2026
671d417
opt(deepseek-v4 projection): modify trace location
lhzhang333 Jun 26, 2026
861600a
update dsv4 flash script
lhzhang333 Jun 26, 2026
a9e9d65
fix(deepseek-v4 projection): validate projection UI controls
lhzhang333 Jun 26, 2026
2c9b4c2
opt(deepseek-v4 projection): sticky model/gpu switcher bar
lhzhang333 Jun 26, 2026
032e87b
gfx1250 mxfp8: skip expert token zero-padding to avoid quant NaN
JohnQinAMD Jun 28, 2026
02b0871
muon: bump emerging_optimizers to 0.4.0a0 + grouped-expert orthogonal…
JohnQinAMD Jun 28, 2026
7137145
v4-pro muon launcher: gfx1250 enablement, gzip traces, compile-disabl…
JohnQinAMD Jun 28, 2026
901facb
feat(muon): batched grouped-expert Newton-Schulz orthogonalize (TP=1)…
JohnQinAMD Jun 28, 2026
6e59a6d
feat(deepseek-v4): fuse the HCA/CSA compressor forward burst
JohnQinAMD Jun 28, 2026
c91a22b
feat(turbo): fuse expert grouped wgrad -> main_grad (beta=1 ACCUMULAT…
JohnQinAMD Jun 28, 2026
c257479
fix(deepseek-v4): forward PRIMUS_TURBO_FUSE_GROUPED_WGRAD into the co…
JohnQinAMD Jun 28, 2026
67a0539
perf(deepseek-v4): architecture-aware tuned defaults for V4 Triton at…
JohnQinAMD Jun 28, 2026
5eea1cb
perf(deepseek-v4): extend arch-aware V4 attn fwd default to the HCA path
JohnQinAMD Jun 29, 2026
24e344e
doc(deepseek-v4): note V4 attn bwd dKV HG default is MQA-only (inert …
JohnQinAMD Jun 29, 2026
3ce6ee4
dsv4 compressed attn launch reduction (#825)
JohnQinAMD Jun 29, 2026
e32f580
fix: fix padding bug and remove use_turbo_permute_padding flag
RuibinCheung Jun 29, 2026
366beb7
chore: remove unused deepseek-v4 launch scripts
wenxie-amd Jun 30, 2026
c069126
feat(deepseek-v4 projection): manual per-layer timing mode
lhzhang333 Jun 30, 2026
f2f2173
feat: update primus turbo wrapper
RuibinCheung Jul 1, 2026
f179a57
feat: add moe_router_padding_for_quantization flag
RuibinCheung Jul 1, 2026
26bd1cc
FlyDSL + Gluon + bugfix + benchmark + refactor (#831)
wenxie-amd Jul 1, 2026
0fe97fd
feat(projection): add DeepSeek-V4 iteration-timeline view + fix fwd/b…
lhzhang333 Jul 2, 2026
1ab1f8e
style: apply pre-commit lint fixes across primus, deepseek-v4, and tests
wenxie-amd Jul 2, 2026
89bae70
perf(v4-attn): optimize `triton_v2` sparse-MLA forward + backward (#838)
wenxie-amd Jul 2, 2026
46509e7
feat(v4-attn): native FlyDSL sparse-MLA backend (flydsl_v1) + triton_…
wenxie-amd Jul 2, 2026
b3d7669
perf(v4): fuse model-body small kernels into Triton (RMSNorm / HC col…
wenxie-amd Jul 3, 2026
9166dd6
[Megatron-LM] fix: duplicated memory footprint when enable turbo grou…
RuibinCheung Jul 6, 2026
1c9a7d9
feat: add clamped swiglu fusion
RuibinCheung Jul 6, 2026
adf496d
feat(v4-attn): gluon_v2 sparse-MLA backend (Gluon fwd + Gluon bwd) (#…
wenxie-amd Jul 6, 2026
538eba4
opt: remove grouped mlp d2h sync
RuibinCheung Jul 7, 2026
4048787
fix: typo and add PRIMUS_BIAS_SWIGLU_FUSION flag for mi455
RuibinCheung Jul 8, 2026
950833e
ci: update Primus-Turbo/AITER pins, disable torch/jax unittests
wenxie-amd Jul 8, 2026
90eebaf
ci: build and install Triton from source in Docker image
wenxie-amd Jul 8, 2026
75cab09
ci: sync ci.yaml + Dockerfiles from main; keep unittest jobs disabled
wenxie-amd Jul 8, 2026
869ea51
opt: add DeepSeek-V4 gluon v3 attention backend (#865)
wenxie-amd Jul 8, 2026
61df353
feat: use triton to replace jit.fuser to avoid torch.compile bug
RuibinCheung Jul 9, 2026
3fd33f1
feat: remove extra d2h and add swiglu fusion for shared experts
RuibinCheung Jul 10, 2026
0ebfa8c
feat: add triton elementwise add kernel for wgrad accumulate
RuibinCheung Jul 10, 2026
e1a541d
feat(deepseek-v4): add Primus-Turbo native-FlyDSL sparse-MLA attentio…
wenxie-amd Jul 14, 2026
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
768 changes: 529 additions & 239 deletions .github/workflows/ci.yaml

Large diffs are not rendered by default.

66 changes: 66 additions & 0 deletions .github/workflows/deploy-projection.yml
Original file line number Diff line number Diff line change
@@ -0,0 +1,66 @@
name: Deploy DeepSeek-V4 Projection

# Publishes the DeepSeek-V4 performance-projection site to GitHub Pages from the
# current development branch (no need to merge to main). The site is served at
# https://<user>.github.io/<repo>/deepseek-v4-projection/
# NOTE: a repo has a single Pages site, so this deployment and the main
# backend-gap dashboard deployment overwrite each other (last run wins). If the
# github-pages environment restricts deployment branches, allow this branch in
# repo Settings -> Environments -> github-pages -> Deployment branches.

on:
workflow_dispatch:
push:
branches:
- dev/tas/deepseek-v4
paths:
- "deepseek-v4/projection/site/**"
- ".github/workflows/deploy-projection.yml"

permissions:
contents: read
pages: write
id-token: write

concurrency:
group: deepseek-v4-projection-pages
cancel-in-progress: true

jobs:
build:
runs-on: ubuntu-latest
steps:
- name: Check out repository
uses: actions/checkout@v4

- name: Assemble Pages bundle
run: |
set -euo pipefail
mkdir -p _site/deepseek-v4-projection
cp -r deepseek-v4/projection/site/. _site/deepseek-v4-projection/
cat > _site/index.html <<'HTML'
<!doctype html>
<meta charset="utf-8" />
<meta http-equiv="refresh" content="0; url=./deepseek-v4-projection/" />
<title>DeepSeek-V4 Projection</title>
<a href="./deepseek-v4-projection/">DeepSeek-V4 Performance Projection</a>
HTML

- name: Configure Pages
uses: actions/configure-pages@v5

- name: Upload Pages artifact
uses: actions/upload-pages-artifact@v3
with:
path: _site

deploy:
needs: build
runs-on: ubuntu-latest
environment:
name: github-pages
url: ${{ steps.deployment.outputs.page_url }}
steps:
- name: Deploy to GitHub Pages
id: deployment
uses: actions/deploy-pages@v4
39 changes: 35 additions & 4 deletions .github/workflows/docker/Dockerfile
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
ARG BASE_IMAGE=docker.io/rocm/primus:v26.2
ARG BASE_IMAGE=docker.io/rocm/primus:v26.3
FROM ${BASE_IMAGE}

ARG PRIMUS_TURBO_COMMIT
Expand All @@ -23,7 +23,7 @@ RUN rm -rf /var/lib/apt/lists/*
ENV ROCSHMEM_HOME=/opt/rocshmem
ENV UCX_HOME=/opt/ucx
# ENV MPI_HOME=/opt/ompi
# Use the system OpenMPI prefix from the v26.2 base image.
# Use the system OpenMPI prefix from the v26.3 base image.
ENV MPI_HOME=/usr/lib/x86_64-linux-gnu/openmpi
ENV ROCM_HOME=/opt/rocm
ENV PRIMUS_TURBO_FRAMEWORK=${PRIMUS_TURBO_FRAMEWORK}
Expand Down Expand Up @@ -55,9 +55,10 @@ RUN rm -rf /opt/rocSHMEM

RUN pip3 uninstall aiter amd-aiter -y || true && \
cd /opt && \
git clone --recursive https://github.com/ROCm/aiter.git && \
git clone https://github.com/ROCm/aiter.git && \
cd aiter && \
git checkout ${PRIMUS_TURBO_AITER_COMMIT} && \
git checkout -f ${PRIMUS_TURBO_AITER_COMMIT} && \
git submodule sync && \
git submodule update --init --recursive && \
PREBUILD_KERNELS=3 pip install --no-cache-dir --use-pep517 .

Expand All @@ -71,6 +72,18 @@ RUN cd /opt && \

RUN rm -rf /opt/Primus-Turbo

# ---------------------------------------------------------------------------
# Install Triton
# ---------------------------------------------------------------------------
RUN cd /opt && \
git clone https://github.com/triton-lang/triton.git && \
cd triton && \
git checkout 09500db9 && \
pip3 install -r python/requirements.txt && \
MAX_JOBS=96 pip3 install --no-build-isolation --force-reinstall --no-deps .

RUN rm -rf /opt/triton

# ---------------------------------------------------------------------------
# Install UCCL-EP (skip for JAX framework)
# ---------------------------------------------------------------------------
Expand All @@ -86,6 +99,24 @@ RUN if [ "$PRIMUS_TURBO_FRAMEWORK" != "JAX" ]; then \
rm -rf /opt/uccl; \
fi

# ---------------------------------------------------------------------------
# Install fixed origami (rocm-libraries@223648a) over the base image's bundled
# 0.1.0. The bundled origami's rank_configs() raises
# `ValueError: vector::reserve` during MoE grouped-gemm kernel selection and
# crashes training (turbo's _safe_rank_configs only catches RuntimeError).
# Skipped for JAX. TODO: drop once a base image ships origami with the fix.
RUN if [ "$PRIMUS_TURBO_FRAMEWORK" != "JAX" ]; then \
rm -rf /tmp/rocm-libraries && \
git clone --filter=blob:none --no-checkout https://github.com/ROCm/rocm-libraries.git /tmp/rocm-libraries && \
cd /tmp/rocm-libraries && \
git sparse-checkout init --cone && \
git sparse-checkout set shared/origami && \
git checkout 223648a26928ebed7f3dd0ccdc044c09f1dccf9b && \
(pip uninstall -y origami || true) && \
pip install --no-cache-dir ./shared/origami/python && \
rm -rf /tmp/rocm-libraries; \
fi

# Set the default working directory
WORKDIR /opt

Expand Down
24 changes: 16 additions & 8 deletions .github/workflows/docker/Dockerfile.ainic
Original file line number Diff line number Diff line change
Expand Up @@ -19,9 +19,9 @@ RUN apt-get update && \
# ---------------------------------------------------------------------------
# Enviroment variables
# ---------------------------------------------------------------------------
ENV WORKDIR=/opt
ENV WORKDIR=/workspace
ENV ROCM_PATH=/opt/rocm
ENV MPI_PATH=/opt/ompi
ENV MPI_PATH=/usr/lib/x86_64-linux-gnu/openmpi

# =============================== Build AINIC Driver ===============================
# WARNING: Please ensure the following environment variables are correctly set:
Expand All @@ -30,16 +30,24 @@ ENV MPI_PATH=/opt/ompi
# WARNING: If these paths are missing, tools and libraries may not function correctly.
# INFO: Installation completed successfully

COPY ${AINIC_BUNDLE_PATH}/ainic_bundle_1.117.5-a-56.tar.gz ${WORKDIR}
COPY ${AINIC_BUNDLE_PATH}/ainic_bundle_1.117.5-a-77.tar.gz ${WORKDIR}
RUN cd ${WORKDIR} && \
rm -rf rccl amd-anp && \
echo "Building ainic bundle... current directory: ${WORKDIR}" && \
tar zxf ainic_bundle_1.117.5-a-56.tar.gz && \
cd ainic_bundle_1.117.5-a-56 && \
tar zxf ainic_bundle_1.117.5-a-77.tar.gz && \
cd ainic_bundle_1.117.5-a-77 && \
tar zxf host_sw_pkg.tar.gz && \
cd host_sw_pkg && \
./install.sh --domain=user -y 2>&1 | tee log_install.txt && \
cd ${WORKDIR} && \
apt-get install -y ./amd/ainic/deb-repo/libionic*.deb
printf '%s\n' \
'Package: libionic1 libionic-dev' \
'Pin: version 54.0-187-1' \
'Pin-Priority: 1001' \
'' \
'Package: perftest' \
'Pin: version 1:25.04.0.0.84-128-1' \
'Pin-Priority: 1001' \
> /etc/apt/preferences.d/ainic-a77 && \
./install.sh --domain=user -y 2>&1 | tee log_install.txt

# ---------------------------------------------------------------------------
# Build rccl
Expand Down
10 changes: 10 additions & 0 deletions .gitignore
Original file line number Diff line number Diff line change
Expand Up @@ -12,3 +12,13 @@ output
experiment
data
pp_simulation_result

# Allow the projection site's breakdown JSON (consumed by the static site /
# GitHub Pages) despite the generic `data` ignore above.
!deepseek-v4/projection/site/data/
!deepseek-v4/projection/site/data/*.json

*.log
*.nohup
*.zip
.triton_cache_shared/
112 changes: 112 additions & 0 deletions check_hc_expand.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,112 @@
#!/usr/bin/env python3
"""Correctness + microbench for the fused HC expand Triton kernel vs eager."""
import importlib.util
import os
import time

import torch

os.environ.setdefault("PRIMUS_HC_EXPAND_TRITON", "1")
# Load the kernel module directly by path to avoid the heavy primus/__init__ chain.
_MOD = os.path.join(
os.path.dirname(os.path.abspath(__file__)),
"primus/backends/megatron/core/transformer/v4_attention_kernels/_triton/hc_expand.py",
)
_spec = importlib.util.spec_from_file_location("hc_expand", _MOD)
_hc = importlib.util.module_from_spec(_spec)
_spec.loader.exec_module(_hc)
hc_expand_triton = _hc.hc_expand_triton


def eager_expand(x, out, post, comb):
write = post.unsqueeze(-1) * out.unsqueeze(-2)
mix = torch.matmul(comb, x)
return write + mix


def run(leading, K, D, dtype):
dev = "cuda"
torch.manual_seed(0)
mk = lambda *s: torch.randn(*s, device=dev, dtype=dtype) # noqa: E731
x = mk(*leading, K, D)
out = mk(*leading, D)
post = mk(*leading, K)
comb = torch.softmax(mk(*leading, K, K).float(), dim=-1).to(dtype)

# fp32 reference (ground truth), plus eager and triton at the test dtype.
x32, o32, p32, c32 = (t.float().clone().requires_grad_(True) for t in (x, out, post, comb))
xa, oa, pa, ca = (t.clone().requires_grad_(True) for t in (x, out, post, comb))
xb, ob, pb, cb = (t.clone().requires_grad_(True) for t in (x, out, post, comb))

ref = eager_expand(x32, o32, p32, c32)
eag = eager_expand(xa, oa, pa, ca)
tri = hc_expand_triton(xb, ob, pb, cb)

g32 = torch.randn_like(ref)
g = g32.to(dtype)
ref.backward(g32)
eag.backward(g)
tri.backward(g)

def maxerr(a, b):
return (a.float() - b.float()).abs().max().item()

tag = f"{tuple(leading)} K={K} D={D} {str(dtype).split('.')[-1]}"
ok = True
for name, r, e, t in [
("fwd", ref, eag, tri),
("dx", x32.grad, xa.grad, xb.grad),
("dout", o32.grad, oa.grad, ob.grad),
("dpost", p32.grad, pa.grad, pb.grad),
("dcomb", c32.grad, ca.grad, cb.grad),
]:
e_err = maxerr(r, e) # eager-vs-fp32 (the dtype's intrinsic noise floor)
t_err = maxerr(r, t) # triton-vs-fp32
# Pass if triton is no worse than eager, with an fp32 reduction-order
# slack relative to the tensor magnitude (length-D sums reorder).
slack = 1e-4 + 5e-5 * r.abs().max().item()
passed = t_err <= max(1.5 * e_err, slack)
ok = ok and passed
print(f" {name:6s} triton_err={t_err:.3e} eager_err={e_err:.3e} {'OK' if passed else 'FAIL'}")
print(f"[{tag}] {'PASS' if ok else 'FAIL'}")
return ok


def bench(leading, K, D, dtype, iters=50):
dev = "cuda"
x = torch.randn(*leading, K, D, device=dev, dtype=dtype, requires_grad=True)
out = torch.randn(*leading, D, device=dev, dtype=dtype, requires_grad=True)
post = torch.randn(*leading, K, device=dev, dtype=dtype, requires_grad=True)
comb = (
torch.softmax(torch.randn(*leading, K, K, device=dev).float(), dim=-1).to(dtype).requires_grad_(True)
)
g = torch.randn(*leading, K, D, device=dev, dtype=dtype)

def step(fn):
for t in (x, out, post, comb):
t.grad = None
y = fn(x, out, post, comb)
y.backward(g)

for fn, name in [(eager_expand, "eager"), (hc_expand_triton, "triton")]:
step(fn)
torch.cuda.synchronize()
t0 = time.perf_counter()
for _ in range(iters):
step(fn)
torch.cuda.synchronize()
ms = (time.perf_counter() - t0) / iters * 1e3
print(f" {name:7s} fwd+bwd {ms:8.3f} ms/iter")


if __name__ == "__main__":
print("=== correctness ===")
all_ok = True
all_ok &= run((4096,), 4, 7168, torch.bfloat16)
all_ok &= run((4096,), 4, 7168, torch.float32)
all_ok &= run((2, 17), 4, 320, torch.bfloat16)
all_ok &= run((3,), 2, 128, torch.float32)
all_ok &= run((5,), 8, 64, torch.bfloat16)
print("\n=== bench (production shape B*S=4096, K=4, D=7168, bf16) ===")
bench((4096,), 4, 7168, torch.bfloat16)
print("\nALL", "PASS" if all_ok else "FAIL")
Loading
Loading