Commit Graph
1418 Commits
Author SHA1 Message Date
Woosuk KwonandCodex 50d812b66a [Perf][DSV4.1] Optimize MegaMoE staging and NVFP4 cache gathers (#57604)
Signed-off-by: Woosuk Kwon <woosuk@inferact.ai>
Co-authored-by: Codex <noreply@openai.com>
2026-09-18 12:57:19 -07:00
frida-andersson 71fc70d3ae [ROCm][DSv4] Fix sparse-indexer logits collapse on gfx950/gfx942 (#50455)
Signed-off-by: Frida Andersson <fanderss@amd.com>
Signed-off-by: frida-andersson <fanderss@amd.com>
2026-09-18 16:57:44 +00:00
Kevin Liandkevinli 911548247a [Distributed][MoonEP] BF16 integration of MoonEP balanced EP backend (#52101)
Signed-off-by: kevinli <306243285+kaijunli-infr@users.noreply.github.com>
Co-authored-by: kevinli <306243285+kaijunli-infr@users.noreply.github.com>
2026-09-18 10:34:01 -04:00
Yuwen ZhouandLi, Jiang <jiang1.li@intel.com> b346479065 [CPU] Add CPU FP8 W8A8 linear/MoE support (#49942)
Signed-off-by: yuwenzho <yuwen.zhou@intel.com>
Co-authored-by: Li, Jiang <jiang1.li@intel.com>
2026-09-18 17:59:42 +08:00
olka-amdandClaude Opus 5 32636580a6 [Bugfix][MLA] TritonMLA: fix illegal memory access on causal multi-token decode (#51065)
Signed-off-by: Oleksandr Kachur <oleksandr.kachur@amd.com>
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
2026-09-18 03:49:53 -05:00
Thien TranandCodex 8da75d6190 C3x SM100 FP8 blockwise - Pad activation scales to multiple of 4 (#57377)
Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg>
Co-authored-by: Codex <noreply@openai.com>
2026-09-18 08:39:27 +00:00
Rohan PotdarandClaude 213c379fca [ROCm][Bugfix] Gate AITER MXFP8 MoE on the aiter enable flag (#57426)
Signed-off-by: Rohan Potdar <rohan.potdar@amd.com>
Co-authored-by: Claude <noreply@anthropic.com>
2026-09-18 00:06:05 +00:00
Misha Goin ac78445bc9 [MoE] Encapsulate TRT-LLM BF16 weight layout handling (#57405)
Signed-off-by: mgoin <mgoin64@gmail.com>
2026-09-17 17:47:30 -04:00
Misha Goin 9a5bd373cf [CI] Retire Weight Loading smoke tests (#57398)
Signed-off-by: mgoin <mgoin64@gmail.com>
2026-09-17 16:04:03 -04:00
Elvir CrnčevićOpenAI Codexmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>Tyler Michael Smith
ac2f0ea82c [MoE][Bugfix] Skip SP padded rows in grouped MoE routing (#56079)
Signed-off-by: Elvir Crncevic <elvircrn@gmail.com>
Co-authored-by: OpenAI Codex <codex@openai.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Co-authored-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
2026-09-17 19:36:19 +00:00
41f9104fa6 [DSv4.1] Integrate Mega-Gate from DeepGEMM (#56266)
Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg>
Co-authored-by: OpenAI Codex <codex@openai.com>
Co-authored-by: zjy0516 <riverclouds.zhu@qq.com>
2026-09-17 12:06:36 -07:00
afriedri a9a7e45f31 [ROCm] Restore VLLM_ROCM_USE_AITER_FP4_ASM_GEMM and default w4a4 ASM GEMM back to off (#57055)
Signed-off-by: Andy Friedrich <afriedri@amd.com>
2026-09-17 13:39:52 -05:00
MattandClaude Opus 5 fef55fd26d [ROCm][CI] Adapt MoE tests to the triton_kernels 3.8 API (#57385)
Signed-off-by: Matthew Wong <Matthew.Wong2@amd.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
2026-09-17 11:21:22 -05:00
Zoe923 6a7de3ba3a [Bugfix][GDN] Fix CuteDSL BF16 KKT inversion divergence (#53864)
Signed-off-by: Zoe923 <871844693@qq.com>
2026-09-17 08:49:56 -07:00
akii96 438434b5b5 [ROCm][Perf] Insert MiniMax-M3 sparse-PA K/V without a contiguous copy (#56849)
Signed-off-by: Aakif Nawaz <aakif.nawaz@amd.com>
2026-09-17 09:44:54 -05:00
78fdf4efb8 [Bugfix][MoE] Convert FlashInfer BF16 weights in place (#54699)
Signed-off-by: Yuchen Wang <93700456+yuchenwang3@users.noreply.github.com>
Signed-off-by: Dakai An <dakaian108@gmail.com>
Co-authored-by: Dakai An <dakaian108@gmail.com>
Co-authored-by: OpenAI Codex <codex@openai.com>
Co-authored-by: aoshen02 <aoshen@inferact.ai>
2026-09-17 22:00:04 +08:00
Rukhaiya2004Akash kaothalkarAntigravityLi, Jiang <jiang1.li@intel.com>
117cf43fb6 [HARDWARE][POWER] Enable W8A8 INT8 MoE on POWER (#55316)
Signed-off-by: Rukhaiya <bibirukhaiya123@gmail.com>
Signed-off-by: Akash kaothalkar <akash.kaothalkar@ibm.com>
Co-authored-by: Akash kaothalkar <akash.kaothalkar@ibm.com>
Co-authored-by: Antigravity <antigravity@google.com>
Co-authored-by: Li, Jiang <jiang1.li@intel.com>
2026-09-17 21:45:13 +08:00
Shantipriya Parida 0eae9acd4d [Bugfix][ROCm][MoE] Fall back instead of crashing when AITER MoE is requested for a non-gated (is_act_and_mul=False) model (#56590)
Signed-off-by: Shantipriya Parida <shantipriya.parida@amd.com>
2026-09-17 13:39:26 +00:00
Lucas WilkinsonandRobert Shaw fe284ea17d [BugFxi] Fix DeepGEMM FP8 workspace over allocation (#53914)
Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>
2026-09-17 09:35:34 -04:00
Simon Danielsson 667b26e50b [Bugfix][ROCm][GLM-5.3-Flash] Apply deferred tilelang.jit already on attribute access (#57192)
Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
2026-09-17 12:02:39 +00:00
08633cb5cd [Qwen3.8-Flash-Next] Enable FP8 TP with FlashInfer TRTLLM MoE (#55867)
Signed-off-by: Canlin <canlinguosdu@gmail.com>
Signed-off-by: Isotr0py <Isotr0py@outlook.com>
Co-authored-by: OpenAI Codex <noreply@openai.com>
Co-authored-by: Isotr0py <Isotr0py@outlook.com>
2026-09-17 08:56:28 +00:00
3e267bae70 [ROCm] triton+triton_kernels 3.8 mxfp4 MoE support (gpt-oss + DeepSeek-V4) (#55934)
Signed-off-by: Rohan Potdar <rohan.potdar@amd.com>
Signed-off-by: Rohan138 <rohanpotdar138@gmail.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Co-authored-by: Matt <156021403+mawong-amd@users.noreply.github.com>
2026-09-16 20:03:34 -05:00
Divakar VermaandAndreas Karatzas 95f4925c3a [ROCm][CI] Enable AITER FP8/unquantized cases in modular-kernel sweep + fix MoRI per-tensor FP8 dispatch (#57074)
Signed-off-by: Divakar Verma <divakar.verma@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
2026-09-16 19:53:06 -05:00
Harry MellorandClaude Opus 5 651a88c09a [CI] Ignore ruff D209, rejoin the docstrings it split, and silence incompatible-rule warnings (#57212)
Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
2026-09-16 10:57:24 -07:00
Harry MellorandClaude Opus 5 8c1557a79c Add pydocstyle to the ruff rules (#52136)
Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
Signed-off-by: Harry Mellor <hej.mellor@gmail.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
2026-09-16 16:44:29 +00:00
d6a1677d55 [Model][DSv4.1] FlashMLA mega attention and the NVFP4 compressed KV cache (#56935)
Signed-off-by: Yongye Zhu <yongye@inferact.ai>
Signed-off-by: 云挚 <ningyunxiao.nyx@antgroup.com>
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Yongye Zhu <yongye@inferact.ai>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-authored-by: Yunxiao Ning <73625538+foraxe@users.noreply.github.com>
Co-authored-by: AI Assistant <noreply@openai.com>
2026-09-16 03:07:50 -07:00
Li, Jiang b68408043d [Bugfix][CPU] Add per-tensor FP8 W8A16 kernel to fix Ministral crash on CPU (#56985)
Signed-off-by: jiang1.li <jiang1.li@intel.com>
2026-09-16 15:49:17 +08:00
ab35354c21 [ROCm][AITER] Skip AITER norm kernels when flattening to 2D would copy (#55991)
Signed-off-by: Zheng Gong <zgong@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: TJian <tunjian.tan@embeddedllm.com>
2026-09-16 07:26:56 +00:00
MattandAndreas Karatzas b3b13c1292 [ROCm][CI] Fix MLA RoPE fused-kernel tests for TheRock image (#57112)
Signed-off-by: Matthew Wong <Matthew.Wong2@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
2026-09-16 05:57:40 +00:00
6ecd97f1b7 [Perf][DSpark] Add KV-only context insertion across V4.1 cache formats (#56441)
Signed-off-by: 0z5a <0z5a@users.noreply.github.com>
Signed-off-by: Yongye Zhu <yongye@inferact.ai>
Co-authored-by: 0z5a <0z5a@users.noreply.github.com>
Co-authored-by: Codex <noreply@openai.com>
Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
2026-09-15 21:34:55 -07:00
maithilijoshi20 22bb158c21 [Bugfix][ROCm] Fix MiniMax-M3 fused MXFP8 block scale (#53674)
Signed-off-by: maithilijoshi20 <97733343+maithilijoshi20@users.noreply.github.com>
2026-09-15 21:59:43 -05:00
ppalangaCursorPoovaiah Palangappamergify[bot] <37929162+mergify[bot]@users.noreply.github.com>Hongxia Yang
a4d2d9d95e [ROCm][Kimi-K3] Enable a4w4 flydsl kernels for KimiK3 (#53940)
Signed-off-by: ppalanga <ppalanga@amd.com>
Signed-off-by: Poovaiah Palangappa <poovaiah.palangappa@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Poovaiah Palangappa <poovaiah.palangappa@amd.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Co-authored-by: Hongxia Yang <62075498+hongxiayang@users.noreply.github.com>
2026-09-16 01:38:02 +00:00
Misha GoinandOpenAI Codex c6fa1f05d1 [Perf][Kernel] Add sampled filtering for persistent top-k (#56346)
Signed-off-by: mgoin <mgoin64@gmail.com>
Co-authored-by: OpenAI Codex <noreply@openai.com>
2026-09-15 14:15:59 -07:00
Jared WenandClaude Fable 5.1 1257512305 [DSA] Wire DeepGEMM sparse MQA logits into the DeepSeek V4.1 indexer (#56254)
Signed-off-by: Jared Wen <w13431838023@gmail.com>
Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com>
2026-09-15 14:07:14 -07:00
Fangzhou Ai dffbb714e4 [ROCm][Perf] Optimize DSV4.1 K=512 decode top-k on gfx950 (#56743)
Signed-off-by: fai <fangzhouai@gmail.com>
2026-09-15 18:34:24 +00:00
35dc273072 [ROCm][Spec Decode] Add Aiter MLA decode support non-causal draft block (#55966)
Signed-off-by: Poovaiah Palangappa <poovaiah.palangappa@amd.com>
Co-authored-by: Poovaiah Palangappa <poovaiah.palangappa@amd.com>
Co-authored-by: John Qin <zhengsheng1997@gmail.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Hongxia Yang <62075498+hongxiayang@users.noreply.github.com>
2026-09-15 12:18:09 -05:00
bb55077411 [Perf][Kernel] Integrate Mega-mHC from DeepGEMM for DeepSeek V4.1 (reopen of #56255) (#56962)
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Signed-off-by: Juntian Liu <Juntianl777@gmail.com>
Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-authored-by: Thien Tran <gau.nernst@yahoo.com.sg>
2026-09-15 04:01:06 -07:00
073f883b56 [Perf][DSpark] Collapse DeepSeek-V4.1 draft states before SP all-gather (#56903)
Signed-off-by: Juntian Liu <Juntianl777@gmail.com>
Co-authored-by: Codex <noreply@openai.com>
Co-authored-by: Woosuk Kwon <woosuk@inferact.ai>
2026-09-15 02:58:57 -07:00
9d14fd3095 [Kernel][MoE] DeepEP v2: async finalize to overlap shared experts with combine (#52781)
Signed-off-by: Mario <404mario@users.noreply.github.com>
Co-authored-by: Mario <404mario@users.noreply.github.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
2026-09-15 00:04:22 -04:00
e6eb0d120c [Model][DSv4.1] Store the whole KV in MXFP8 (FlashMLA V4.1 record) (#56893)
Signed-off-by: Yongye Zhu <yongye@inferact.ai>
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Yongye Zhu <yongye@inferact.ai>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
2026-09-14 20:12:55 -07:00
JohnQinAMDandClaude Opus 5 ef5f7cd119 [ROCm][DSv4.1][Perf] Dequantize the MXFP8 weight once when dot_scaled cannot be used (#56560)
Signed-off-by: JohnQinAMD <yanyuan.qin@amd.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
2026-09-15 09:07:26 +08:00
Fangzhou AiandCursor 00972dfd72 [ROCm][DSV4.1][Perf] Fold the mHC post step into the delayed pre projection (#56513)
Signed-off-by: Fangzhou Ai <fangzhou@semianalysis.ai>
Signed-off-by: Fangzhou Ai <fangzhou.ai@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
2026-09-14 16:44:56 -05:00
Fangzhou AiandCursor dabc4362b4 [ROCm][DSV4.1][Perf] Stride the DSA decode candidate mask over the live context (#56628)
Signed-off-by: Fangzhou Ai <fangzhou.ai@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
2026-09-14 20:47:32 +00:00
sychen52 7702ee87db [Bugfix][DSA] Write nvfp4_ds_mla from the fused norm+rope kernel (#55538)
Signed-off-by: Shiyang Chen <shiychen@nvidia.com>
2026-09-15 00:25:33 +04:00
Netanel Haber d4ee7fe7a9 [Quantization] Select linear backends per quantization (#51204)
Signed-off-by: Netanel Haber <58652339+netanel-haber@users.noreply.github.com>
2026-09-14 16:24:18 -04:00
Yongye Zhu 5372e72a98 [Perf][DSv4.1] Fold the mHC post block into the delayed pre projection (#56633)
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
2026-09-14 18:57:37 +00:00
Samuel NordmannmgoinClaude Opus 4.8ElizaWszolaOpenAI Codexmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
03dc26e639 [Perf][Kernel][Quantization] Fuse ReLU2 with static FP8 activation quantization (#53793)
Signed-off-by: mgoin <mgoin64@gmail.com>
Signed-off-by: ElizaWszola <ewszola@redhat.com>
Signed-off-by: snordmann <snordmann@nvidia.com>
Co-authored-by: mgoin <mgoin64@gmail.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Co-authored-by: ElizaWszola <ewszola@redhat.com>
Co-authored-by: OpenAI Codex <noreply@openai.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-09-14 10:26:15 -07:00
Matthias Gehre dc2e8f1157 [ROCm][Perf] W4A16: keep skinny GEMM zero-points packed 4-bit (#54965)
Signed-off-by: Matthias Gehre <matthias.gehre@amd.com>
2026-09-14 17:40:25 +08:00
78e84261ab [Performance] Add Triton kernel for Gemma3n sparse GELU (#48498)
Signed-off-by: BWAAEEEK <jooho414@gmail.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: Cyrus Leung <tlleungac@connect.ust.hk>
2026-09-14 07:42:15 +00:00
Bugen Zhao cf1584f373 [Refactor] Normalize DeepSeek V4.1 model package naming (#56741)
Signed-off-by: Bugen Zhao <i@bugenzhao.com>
2026-09-13 23:17:50 -07:00