Commit Graph
1210 Commits
Author SHA1 Message Date
Guanxin LiandClaude Opus 5 925ea7e60f [CI][Test] Seed the DeepEP v2 MoE workers, not just the parent (#50589)
Signed-off-by: Guanxin Li <guanxinl@nvidia.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
2026-08-14 16:01:24 -04:00
Zhuobin Huangandnodeeeeee 63a9a5010a [Attention][DSA] Take the native decode path for MTP=3 on SM90 (#52164)
Signed-off-by: zobinHuang <zobin1999@gmail.com>
Co-authored-by: nodeeeeee <zhangkai.nodeee@gmail.com>
2026-08-14 10:09:51 +00:00
ccaadaroandLi, Jiang <jiang1.li@intel.com> bda4c3e8ee [CPU] Fold the MXFP4 block scale in 2 instructions instead of 4 (#51583)
Signed-off-by: ccaadaro <ccanadar@unex.es>
Co-authored-by: Li, Jiang <jiang1.li@intel.com>
2026-08-14 08:19:21 +00:00
1be3628367 [Kernel][Perf] Add fused CUDA post-conv MTP decode kernel for Qwen3.5 GDN (#51674)
Signed-off-by: Jie Fang <jief@nvidia.com>
Co-authored-by: OpenAI Codex <codex@openai.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
2026-08-14 11:44:45 +08:00
Luke Alonso 3c79b1a8bf [Kernel] Add B12X dense linear backends (#52016)
Signed-off-by: Luke Alonso <lalonso@gmail.com>
2026-08-13 19:48:33 -07:00
bastefaniak 71b0da7c4e [Bugfix] Fix .../mrope.py::apply_interleaved_rope() when torch.compile is used in torch==2.13 (#52005)
Signed-off-by: Bartosz Stefaniak <bstefaniak@nvidia.com>
2026-08-13 22:08:40 +00:00
stefankoncarevic 11c3fa4adc [Bugfix][ROCm][CI] Give the AITER MLA decode metadata stub its MLA dims (#52139)
Signed-off-by: Stefan Koncarevic <Stefan.Koncarevic@amd.com>
2026-08-13 13:35:00 -05:00
fxmarty-amdAndreas Karatzasmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
c5b7c069a4 [ROCm] Defer tilelang import through its import from vllm.tilelang_utils import tilelang and relaxed has_tilelang (#51159)
Signed-off-by: Felix Marty <Felix.Marty@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-08-13 13:12:40 -05:00
Thien Tran 6014f9e672 [Kimi-K3] Add GEMM-RS for sequence parallelism (#52079)
Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg>
2026-08-13 09:02:28 -07:00
Michael Goin 2d24355eb8 [Bugfix] Fix packed GDN decode launch for large batch-head grids (#52030)
Signed-off-by: mgoin <mgoin64@gmail.com>
2026-08-13 07:19:08 -07:00
Yongye ZhuClaude Opus 5mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
903d2efe7e [Attention][MLA] Fuse Kimi-K3 chunked-context K/V packing (#51772)
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-08-12 22:42:57 -07:00
61826c1c6f [Hardware][Power] Unqualized MoE Backend for Power (VSX) (#51624)
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>
2026-08-13 10:56:09 +08:00
xiao fengandMichael Goin e62abc37d4 [Bugfix][Kernel] Fix persistent top-k histogram reuse after short rows (#49139)
Signed-off-by: fxfxfxfxfxfxfxfx <227935476@qq.com>
Co-authored-by: Michael Goin <mgoin64@gmail.com>
2026-08-12 12:54:45 -07:00
kiroxuCopilotWentao Yemergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
19d61b1273 [Bugfix] Initialize DeepGemmQuantScaleFMT oracle lazily; bound QuantFP8 UE8M0 packed path to group_size 128 (#51359)
Signed-off-by: BabyDrangoner <148877251+BabyDrangoner@users.noreply.github.com>
Signed-off-by: kiroxu <148877251+BabyDrangoner@users.noreply.github.com>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Co-authored-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-08-12 14:27:50 -04:00
stefankoncarevic d20a031d33 [Bugfix][ROCm][MoE] Update AITER MXFP4 W4A16 tests to the renamed expert_mask (#51980)
Signed-off-by: Stefan Koncarevic <Stefan.Koncarevic@amd.com>
2026-08-12 17:45:45 +00:00
9035151d6c [Model] Add native Dots3 NOTE multimodal support (#51255)
Signed-off-by: KurodaKanbei <mistergalahad@gmail.com>
Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>
Co-authored-by: OpenAI Codex <codex@openai.com>
Co-authored-by: Jee Jee Li <jeejeelee@inferact.ai>
Co-authored-by: Jee Jee Li <pandaleefree@gmail.com>
2026-08-12 23:13:20 +08:00
Matvei PashkovskiiandClaude 20727e2841 [Hardware][AMD] Enable fused bf16→fp32 router GEMM on ROCm (#50268)
Signed-off-by: Matvei Pashkovskii <mpashkov@amd.com>
Co-authored-by: Claude <noreply@anthropic.com>
2026-08-12 02:46:13 +00:00
Micah WilliamsonandAndreas Karatzas 045a986f3b [ROCm][CI] Speed Up ROCm Skinny GEMM Tests (reduced parameterizations, (#51877)
Signed-off-by: Micah Williamson <micah.williamson@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
2026-08-12 01:40:13 +00:00
1d2d83a07f [Attention] Add FlashInfer XQA decode support on SM12x (#49718)
Signed-off-by: Andrii Skliar <askliar@nvidia.com>
Signed-off-by: Andrii Skliar <andreyws96@gmail.com>
Co-authored-by: Andrii Skliar <askliar@nvidia.com>
Co-authored-by: OpenAI Codex <codex@openai.com>
2026-08-11 16:32:49 -07:00
Rohan PotdarandClaude Opus 4.8 0f0cb918b7 [ROCm][MoE] Fix expert_map vs AITER expert_mask for non-AITER experts under EP (#49758)
Signed-off-by: Rohan Potdar <rohan.potdar@amd.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
2026-08-11 15:33:02 -05:00
Andreas Karatzasandmergify[bot] <37929162+mergify[bot]@users.noreply.github.com> 36f4630d89 [Bugfix][MoE] Fix fused block-scale orientation (#50727)
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-08-11 16:05:38 -04:00
Andreas KaratzasandOpenAI Codex 1ab2801dde [ROCm] Remove stale SDPA and skinny GEMM workarounds (#50907)
Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com>
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: OpenAI Codex <codex@openai.com>
2026-08-11 14:19:16 -04:00
aditi-amdrootAndreas Karatzasmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
5426311d91 [Kernel][ROCm][Perf] FlyDSL decode-attention kernel for 4-bit TurboQuant KV cache (#47896)
Signed-off-by: root <root@smci355-ccs-aus-m02-09.cs-aus.dcgpu>
Co-authored-by: root <root@smci355-ccs-aus-m02-09.cs-aus.dcgpu>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-08-11 08:45:39 -07:00
pmanczakandKunshang Ji 5b5eae2394 [Misc] Enable test_silu_mul_fp8_quant_deep_gemm on XPU (#49444)
Signed-off-by: pmanczak <pawel.manczak@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-08-11 23:19:14 +08:00
Fangzhou Ai 1044d19998 [ROCm][DSV4] Preserve native MXFP4 TP8 shard allocation (#51473)
Signed-off-by: Fangzhou Ai <fangzhou.ai@amd.com>
2026-08-10 23:57:27 -07:00
Priyansh JainandAndreas Karatzas 2acb055ed7 [CPU][Zen] Route BF16 MoE inference through zentorch on AMD (#44201)
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
2026-08-10 23:20:05 -07:00
c76a425278 [Kernel] Optimize long-context MLA cache gathers (#51739)
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: OpenAI Codex <noreply@openai.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
2026-08-10 22:23:43 -07:00
63ac04a61e [Kimi-K3] DCP support (#50484)
Signed-off-by: girasoley <girasoley@inferact.ai>
Signed-off-by: Summer Yang <girasoleyang@gmail.com>
Signed-off-by: 云挚 <ningyunxiao.nyx@antgroup.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: 云挚 <ningyunxiao.nyx@antgroup.com>
Co-authored-by: foraxe <73625538+foraxe@users.noreply.github.com>
Co-authored-by: Woosuk Kwon <woosuk@inferact.ai>
2026-08-10 12:54:26 -07:00
fanxingran 37fbf52084 [ROCm][MLA] [K3] Fix fp8 KV cache decode on the AITER MLA backend (#51011)
Signed-off-by: fanxingran <xingran.fan@amd.com>
2026-08-10 11:42:10 -04:00
Francesco FuscoandClaude fac808b36f [Perf][Hybrid] 3D-grid tiling of the state-copy Triton kernels (#49436)
Signed-off-by: Francesco Fusco <ffu@zurich.ibm.com>
Co-authored-by: Claude <noreply@anthropic.com>
2026-08-10 17:02:49 +02:00
pmanczak ec21f61d82 [Misc] Enable test_fused_moe_wn16 on XPU (#51672)
Signed-off-by: pmanczak <pawel.manczak@intel.com>
2026-08-10 21:57:02 +08:00
Andreas KaratzasandMicah Williamson 31cd109f18 [ROCm][CI] Extend ROCm AITER MHA (FA) coverage (#40958)
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Signed-off-by: Micah Williamson <micah.williamson@amd.com>
Co-authored-by: Micah Williamson <micah.williamson@amd.com>
2026-08-10 15:35:13 +08:00
Lena OnyshchenkoandKunshang Ji 751f2ccdd3 [Kernel][XPU] Tensor-descriptor operand loads for Triton W8A8 scaled_mm (#47205)
Signed-off-by: oonyshch <xonyshch@gmail.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-08-10 00:56:11 +00:00
Isotr0py 073c510c91 [Migration] Migrate bitsandbytes support to OOT plugin (#43529)
Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn>
Signed-off-by: Isotr0py <Isotr0py@outlook.com>
2026-08-09 10:27:31 +08:00
Aarushi JainandAarushi Jain 7581c56c86 [Test] Add ROCm AITER FP8 MLA prefill accuracy test (#51457)
Signed-off-by: Aarushi Jain <aarushi@amd.com>
Co-authored-by: Aarushi Jain <aarushi@amd.com>
2026-08-08 19:42:26 -05:00
Tianmu Li 700d39b558 [CPU] Optimize routed FP8/MXFP4 MoE GEMM dispatch (#50949)
Signed-off-by: Li, Tianmu <tianmu.li@intel.com>
2026-08-08 12:56:14 +08:00
Wei-Ming Chen 1c94e8dc7d Add NVFP4 KV 4-over-6 scale search (#45187)
Signed-off-by: weimingc <17592131+meenchen@users.noreply.github.com>
2026-08-07 13:22:27 -07:00
Walter Beller-MoralesandHarry Mellor 70456e5e6f [Bugfix] when loading weights skip empty expert bias if model does not support them (#50937)
Signed-off-by: walterbm <walter.beller.morales@gmail.com>
Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
2026-08-07 12:45:52 -07:00
fxmarty-amd 3518110f2f [Online quantization] Add online MXFP4 quantization support (#49347)
Signed-off-by: Felix Marty <Felix.Marty@amd.com>
2026-08-07 12:39:46 -07:00
Aarushi JainandAarushi Jain a0056e103e [Test] Add ROCm AITER MLA op registration and env gating tests (#50930)
Signed-off-by: Aarushi Jain <aarushi@amd.com>
Co-authored-by: Aarushi Jain <aarushi@amd.com>
2026-08-07 14:07:31 -05:00
Jinzhen LinandOpenAI Codex fcde8e1460 [Refactor] refactor humming linear and moe backends to use explicit layer configs (#49610)
Signed-off-by: Jinzhen Lin <jinzhen.ljz@antgroup.com>
Co-authored-by: OpenAI Codex <noreply@openai.com>
2026-08-07 10:03:19 -07:00
Julien Debache ae934ba8a5 feat: extended EPLB support for Mistral Large 3 and additional MoE backends (#48355)
Signed-off-by: jdebache <jdebache@nvidia.com>
2026-08-07 05:32:52 -07:00
pmanczakandpmanczak 6b5bec7bed [Misc] Add and enable Triton kernel unit tests on XPU (#45694)
Signed-off-by: pmanczak <pmanczak@users.noreply.github.com>
Signed-off-by: pmanczak <pawel.manczak@intel.com>
Co-authored-by: pmanczak <pmanczak@users.noreply.github.com>
2026-08-07 02:53:03 -07:00
stefankoncarevicAndreas Karatzasmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
0de0362ea1 [ROCm][CI] Loosen block-FP8 fused MoE test tolerance for large-K shapes (#48847)
Signed-off-by: Stefan Koncarevic <Stefan.Koncarevic@amd.com>
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-08-07 01:19:26 -07:00
e08111211b [Kernel] Support Nvfp4 Cutedsl Moe Swiglu-oai and Relu2(non-gated) Activation (#47106)
Signed-off-by: Tiekai Bi <tiekaib@nvidia.com>
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>
2026-08-06 17:42:45 -07:00
d8eabdbfbe [ROCm][MLA] Use asm decode for non-divisor small head counts (#50578)
Signed-off-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>
Signed-off-by: Markus Hartikainen <markus.hartikainen@amd.com>
Signed-off-by: Liuyinfeng01 <yinfeliu@amd.com>
Co-authored-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Markus Hartikainen <markus.hartikainen@amd.com>
Co-authored-by: Liuyinfeng01 <yinfeliu@amd.com>
2026-08-06 14:56:53 -04:00
Matthew BonanniClaude Opus 5OpenAI Codexmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
b38e111d3e [Attention][MLA] Per-request scheduling for MLA chunked context (#50613)
Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-authored-by: OpenAI Codex <noreply@openai.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-08-06 16:45:44 +00:00
zofiaandKunshang Ji 8060260e1d [Linear] [Kernel] add block-wise scaled_mm (#49932)
Signed-off-by: Zhu, Zufang <zufang.zhu@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-08-06 14:40:26 +08:00
Gabriel WuandYongye Zhu 2e35c529b8 [Bugfix][Quantization] Fix MXFP4 conversion for FlashInfer CUTLASS (#51038)
Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>
2026-08-05 23:19:22 -07:00
Aarushi Jain 71f975af43 [ROCm][CI] Add MLA decode accuracy and determinism tests (#50480)
Signed-off-by: aarushjain29 <Aarushi.Jain2@amd.com>
2026-08-05 20:35:56 -05:00