Commit Graph
938 Commits
Author SHA1 Message Date
Fadi ArafehGitHubLi, Jiang <jiang1.li@intel.com>
0b25cf4419 [CPU][Perf] Enable fused kernels for GDN's gated delta rules (#43534)
Signed-off-by: Fadi Arafeh <fadi.arafeh@arm.com>
Co-authored-by: Li, Jiang <jiang1.li@intel.com>
2026-06-02 08:00:48 +00:00
559d6710bf [PERF]MiniMax-M2 gate kernel (#38445)
Signed-off-by: Jee Jee Li <pandaleefree@gmail.com>
Signed-off-by: qianlihuang <91178480+qianlihuang@users.noreply.github.com>
Co-authored-by: Yiliu Dong <91178480+qianlihuang@users.noreply.github.com>
2026-05-29 18:28:34 -07:00
bnellnmandGitHub 7b98f498cd [MoE Refactor] Remove supports_expert_map (#43108)
Signed-off-by: Bill Nell <bnell@redhat.com>
2026-05-29 17:26:56 -04:00
JartXandGitHub 0cff0741ff [Kernel][ROCm] Native W4A16 kernel for AMD RDNA3 (gfx1100) — fp16 + bf16 (#41394)
Signed-off-by: JartX <sagformas@epdcenter.es>
2026-05-29 11:04:40 +00:00
Woosuk KwonandGitHub 7bd45da585 [DSv4] Move mHC tilelang kernels & Don't use CustomOP in dsv4/nvidia (#43905)
Signed-off-by: Woosuk Kwon <woosuk@inferact.ai>
2026-05-29 10:25:02 +08:00
f3b2a819f7 [Perf][KDA] Fuse gate softplus, chunk-local cumsum, and RCP_LN2 scaling (#43667)
Signed-off-by: haojiangzheng <justineric096@gmail.com>
Co-authored-by: haojiangzheng <justineric096@gmail.com>
2026-05-28 13:47:08 +00:00
Wentao YeandGitHub 64e1218673 [Perf] Optimize moe permute by pre-allocate buffer, 9~14% kernel performance improvement (#43014)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
2026-05-28 06:18:26 -07:00
TJianandGitHub 0ba46d4b11 [ROCm][DSV4] Enable Tilelang MHC replacing torch/triton mhc (#43679)
Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com>
2026-05-28 07:05:28 +00:00
Thien TranandGitHub e54eff769d [Bugfix] Pass routed_scaling_factor to FlashInfer TRTLLM BF16 MoE (#43769) 2026-05-27 21:29:14 -07:00
Yongye ZhuandGitHub 2d2c660104 [MoE] Remove inplace fused experts mechanism (#43727)
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
2026-05-27 20:00:19 -07:00
jatseng-aiandGitHub 05c50c721e [ROCm] mori: add InterNodeV1LL inter-node kernel selection via VLLM_MORI_INTERNODE_KERNEL (#41751)
Signed-off-by: jatseng-ai <jatseng@amd.com>
2026-05-28 00:33:32 +08:00
Jee Jee LiandGitHub 6e503868ca [Kernel] Porting fuse_minimax_qk_norm to manual fusion (#43410)
Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>
2026-05-26 13:16:03 -07:00
f51bbc694d [MoE Refactor] W4a8 int8 oracle (#42789)
Signed-off-by: Bill Nell <bnell@redhat.com>
Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>
2026-05-26 11:15:42 -04:00
Yongye ZhuandGitHub 6ab6ffb428 [Feat][DSV4] Fuse q pad into deepseek v4 fused kernel (#43162) 2026-05-26 05:12:54 -10:00
Thien TranandGitHub d56612c621 [GDN] GDN Prefill kernel for SM100 (#43273)
Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg>
2026-05-26 14:02:11 +08:00
d56285c747 Tuning script and configs for Triton Mamba SSU kernel (#43083)
Signed-off-by: Banani Ghosh <bg2502@nyu.edu>
Signed-off-by: Daniel Serebrenik <daserebrenik@nvidia.com>
Co-authored-by: Banani Ghosh <bg2502@nyu.edu>
2026-05-24 20:12:44 +03:00
Andreas KaratzasandGitHub d28bdf9344 [ROCm][CI] Fix ROCm LoRA Transformers fallback with full CUDA graphs (#41577)
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
2026-05-23 04:31:32 +00:00
gnovackandGitHub f743254143 DSv4 fused Q-norm kernel grid refactor (#42353) 2026-05-22 15:21:33 -07:00
Yongye ZhuandGitHub 843715739b [Refactor] Extract DeepSeek V4 sparse MLA impl into model folder (#43149) 2026-05-22 10:06:31 -07:00
sychen52andGitHub fb21d8b4f9 Add NVFP4 MOE support for Deepseek V4. (#42209)
Signed-off-by: Shiyang Chen <shiychen@nvidia.com>
2026-05-22 07:21:51 -07:00
68e07d5916 [Bug] Fix ci issue assert output_size is not None AssertionError (#43261)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Signed-off-by: Isotr0py <Isotr0py@outlook.com>
Co-authored-by: Isotr0py <Isotr0py@outlook.com>
2026-05-21 16:58:09 +08:00
Michael GoinandGitHub 5774aad9c5 [Perf][gpt-oss] Downgrade triton_kernels to v3.5.1 (#43135)
Signed-off-by: mgoin <mgoin64@gmail.com>
2026-05-20 14:13:12 -07:00
363fc84407 Integrate flashinfer b12x MoE and FP4 GEMM kernels for SM120/121 (#40082)
Signed-off-by: Meenakshi Venkataraman <meenakshiv@nvidia.com>
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
2026-05-20 17:21:11 +00:00
Dom BrownandGitHub d247a931cc [feat] Add FP8 per-tensor Q scale support to Triton attention backend (#42080)
Signed-off-by: Dom Brown <3886319+DomBrown@users.noreply.github.com>
2026-05-19 09:02:05 -07:00
Woosuk KwonandGitHub b14be81c1f [Model Refactoring] Move deepseek_v4_ops to models/deepseek_v4 [3/N] (#43073)
Signed-off-by: Woosuk Kwon <woosuk@inferact.ai>
2026-05-19 00:52:54 -07:00
3ca8db2ef8 add cutedsl dsv4 indexer fp8 kernel (#42899)
Signed-off-by: george <george@inferact.ai>
Co-authored-by: george <george@inferact.ai>
2026-05-18 21:17:56 -07:00
Wentao YeandGitHub 37ece593c1 [Perf] Padded nvfp4 quant kernel to remove additional copy, 2.4%~5.7% e2e performance improvement (#42774)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
2026-05-18 16:38:12 -07:00
Siddharth BedekarGitHubRobert ShawOpenAI Codexgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
78e7a7b9b0 Refactor AWQ Marlin MoE onto modular WNA16 oracle (#42483)
Signed-off-by: Siddharth Bedekar <bedeksid@gmail.com>
Signed-off-by: Siddharth Bedekar <104613085+bedeks@users.noreply.github.com>
Co-authored-by: Robert Shaw <robertgshaw2-redhat@users.noreply.github.com>
Co-authored-by: OpenAI Codex <codex@openai.com>
Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-05-18 08:02:43 -07:00
Yuwen ZhouandGitHub 88a860d754 [CPU] Add MXFP4 W4A16 MoE support (#41922)
Signed-off-by: yuwenzho <yuwen.zhou@intel.com>
Signed-off-by: Yuwen Zhou <yuwen.zhou@intel.com>
2026-05-18 03:04:45 -07:00
0867497368 [CI/Build] Bump flashinfer to v0.6.11.post2 (#41711)
Signed-off-by: Artem Perevedentsev <aperevedents@nvidia.com>
Co-authored-by: Vadim Gimpelson <156319763+vadiklyutiy@users.noreply.github.com>
2026-05-16 14:55:12 -07:00
b2c58ee942 [FlashAttn] Fix supports_kv_cache_dtype() accepting unhandled fp8 kv-cache dtype variants (#42685)
Signed-off-by: Lanze Liu <lanzetech@gmail.com>
Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>
Co-authored-by: Matthew Bonanni <mbonanni@redhat.com>
2026-05-15 15:34:59 -04:00
06d020bb6e [Bugfix] Fix SM121 (DGX Spark) exclusion from Marlin/CUTLASS FP8 paths (#35568)
Signed-off-by: Blake Ledden <blake@secondnaturecomputing.com>
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
Co-authored-by: Pavani Majety <pmajety@nvidia.com>
2026-05-15 10:59:00 -07:00
f351455f0f [CPU][RISC-V] Add RVV-optimized attention kernels for RISC-V Vector Extension (#40119)
Signed-off-by: liuyudong <liuyudong@iscas.ac.cn>
Co-authored-by: Claude <noreply@anthropic.com>
2026-05-15 12:08:23 +08:00
f8848b2f2d [Bugfix] Add swiglu limits to deepgemm fp8 methods (#41986)
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>
2026-05-14 15:43:13 -07:00
a8887c208f [Bugfix] [ROCm] [DSV4] [Perf] Add aiter mhc support (#41946)
Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com>
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>
2026-05-13 21:43:15 +08:00
3c413a5481 Triton attention: add USE_TD constexpr for tensor descriptor Q/K/V load/store (#40327)
Signed-off-by: Artur Fierka <artur.fierka@intel.com>
Co-authored-by: quinnlp <quinnlp@users.noreply.github.com>
2026-05-13 13:57:41 +02:00
bnellnmandGitHub 18f6bf5a21 [MoE Refactor] Add sequence parallel tests to test_moe_layer.py (#41299)
Signed-off-by: Bill Nell <bnell@redhat.com>
2026-05-12 21:52:19 -04:00
d9b4990783 [MoE Refactor] EPLB refactoring for FusedMoE (#41055)
Signed-off-by: Bill Nell <bnell@redhat.com>
Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>
2026-05-12 14:16:31 -04:00
4d591db470 [MoE Refactor] Introduce RoutedExperts alias for FusedMoE and don't store SharedExperts in MK (#40735)
Signed-off-by: Bill Nell <bnell@redhat.com>
Signed-off-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>
Signed-off-by: Robert Shaw <robertgshaw2@gmail.com>
Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>
Co-authored-by: Robert Shaw <robertgshaw2@gmail.com>
2026-05-12 13:37:44 -04:00
bnellnmandGitHub 6427603ae8 [MoE Refactor] Move remaining experts classes to experts directory (#42334)
Signed-off-by: Bill Nell <bnell@redhat.com>
2026-05-12 09:19:46 -04:00
206eaed08d [MoE Refactor] Move expert map related code into ExpertMapManager class (#41046)
Signed-off-by: Bill Nell <bnell@redhat.com>
Signed-off-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>
Signed-off-by: Robert Shaw <robertgshaw2@gmail.com>
Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>
Co-authored-by: Robert Shaw <robertgshaw2@gmail.com>
2026-05-12 09:18:27 -04:00
Hexiang WangandGitHub 7863fff6e5 [ROCm][DSv4] implement flash sparse mla with triton kernels (#41812)
Signed-off-by: whx-sjtu <xiaowang990929@gmail.com>
2026-05-11 09:27:11 -07:00
724ed2fc35 [DSv4] Improved dequant gather K cache kernel (#42236)
Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg>
Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>
2026-05-11 10:41:12 -04:00
wangxiyuanandGitHub 5536fc0c01 [Misc] Replace mamba_type string literals with MambaAttentionBackendEnum (#41188)
Signed-off-by: wangxiyuan <wangxiyuan1007@gmail.com>
2026-05-11 03:59:36 +00:00
171019ab19 add fused mhc_post_pre kernel (#41536)
Signed-off-by: george <george@inferact.ai>
Co-authored-by: george <george@inferact.ai>
2026-05-10 19:56:52 -07:00
1b57eb41f2 [MoE] Move various experts classes to fused_moe/experts/ (#41979)
Signed-off-by: Jackmin801 <ongjackm@gmail.com>
Signed-off-by: Robert Shaw <robertgshaw2@gmail.com>
Signed-off-by: Jackmin801 <56836461+Jackmin801@users.noreply.github.com>
Signed-off-by: Bill Nell <bnell@redhat.com>
Co-authored-by: Jackmin801 <ongjackm@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: Robert Shaw <robertgshaw2@gmail.com>
Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>
Co-authored-by: Jackmin801 <56836461+Jackmin801@users.noreply.github.com>
2026-05-11 07:54:33 +08:00
Thien TranandGitHub 530d371302 [DSv4] Improved fused Indexer Q quant kernel (#41428) 2026-05-09 01:20:32 -07:00
Michael GoinandGitHub 8bcd8a260c [Bugfix] Fix FlashInfer CUTLASS MXFP4-MXFP8 MoE by restoring swizzled scale (#42089)
Signed-off-by: mgoin <mgoin64@gmail.com>
2026-05-08 15:59:06 -07:00
John CalderonandGitHub 8a2fc80b84 [CUDA][CUTLASS] Enable cutlass scaled mm for non-compatible sizes (#41868)
Signed-off-by: John Calderon <jcalderon@nvidia.com>
2026-05-08 15:58:05 -07:00
0b99971352 [Kernel][Helion] Optimize Helion config parsing latency (#40850)
Signed-off-by: Yanan Cao <gmagogsfm@gmail.com>
Co-authored-by: Claude Sonnet 4 <noreply@anthropic.com>
2026-05-07 20:27:34 -07:00