Commit Graph
71 Commits
Author SHA1 Message Date
Divakar Verma 166d14e9bf [bugfix] skip conch kernel for g_idx reordering (#45072)
Signed-off-by: Divakar Verma <divakar.verma@amd.com>
2026-06-10 23:04:19 +08:00
Michael GoinandClaude 87c15d46e3 [Bugfix] Lazily import the humming quantization backend (#44921)
Signed-off-by: mgoin <mgoin64@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>
2026-06-10 06:06:17 -07:00
Michael Goin 6afa25000c [Bugfix] Canonicalize FP8 weight layout to (K, N) at the source (#44735)
Signed-off-by: mgoin <mgoin64@gmail.com>
2026-06-08 14:37:36 -06:00
Ma Jian eebce65756 [XPU]feat: add DeepSeek-V4 XPU attention decode path (#42953)
Signed-off-by: Ma Jian <jian1.ma@intel.com>
2026-06-08 13:27:12 +08:00
Yongye ZhuandClaude ec0a31d4aa [Bugfix][Kernel] Fix mHC fused-RMSNorm big-fuse miscompile for hidden_size != 4096 (#44692)
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>
2026-06-06 10:44:21 +08:00
Han Lingemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>TJian
165b7864d0 [ROCM] [FEAT] Integrate Aiter hipBLASLt GEMM online tuning (#40426)
Signed-off-by: hanlin12 <hanlin12@amd.com>
Signed-off-by: Han Lin <hanlin12@amd.com>
Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
Co-authored-by: TJian <tunjian.tan@embeddedllm.com>
2026-06-04 23:45:36 -07:00
Michael Goin 06ee2d8433 [Quant] Support compressed-tensors WNA8O8Int linears and WNInt embeddings (#44340)
Signed-off-by: mgoin <mgoin64@gmail.com>
2026-06-04 07:40:33 -07:00
Dipika Sikka e68988a248 Refactor CT NVFP4 linear to use a single class (#42443) 2026-06-04 08:25:08 -04:00
e5232679a3 [XPU] Add XPU block-scaled W8A8 fp8 path (#39968)
Signed-off-by: Wu, Xiaochang <xiaochang.wu@intel.com>
Signed-off-by: Xiaochang Wu <xiaochang.wu@intel.com>
Co-authored-by: Yuxiang <yuxiang.liang@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-06-03 20:16:19 +08:00
zofiaandKunshang Ji 6314de8bad [XPU] [Bug] remove xpuw4a16 output size check (#44168)
Signed-off-by: Zhu, Zufang <zufang.zhu@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-06-02 22:26:20 +08:00
Wentao Yeandmergify[bot] <37929162+mergify[bot]@users.noreply.github.com> 985c97a6a8 [Perf] Optimize cutlass fp8 scaled mm bypassing padding, 20% kernel performance improvement (#43706)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-06-01 09:05:21 -04:00
3fd9d2d357 [CPU][Zen] Route W8A8 and W4A16 linear inference through zentorch on AMD Zen CPUs (#41813)
Signed-off-by: R <Ganesh.R@amd.com>
Signed-off-by: Harshal Adhav <harshal.adhav@amd.com>
Signed-off-by: Aakar Dwivedi <aadwived@amd.com>
Co-authored-by: R <Ganesh.R@amd.com>
Co-authored-by: Harshal Adhav <harshal.adhav@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Michael Goin <mgoin64@gmail.com>
2026-05-30 14:17:21 -05:00
JartX 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 Kwon 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
Li, Jiang 20d69d100a [CPU] Migrate cpu_awq into awq_marlin (#43841)
Signed-off-by: jiang1.li <jiang1.li@intel.com>
2026-05-28 22:36:31 +08:00
TJian 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
Luciano MartinsandLuciano Martins 2272062471 [Kernel] Enable TritonW4A16LinearKernel as CUDA fallback for non-Marlin-aligned W4A16 shapes (#43731)
Signed-off-by: Luciano Martins <lucianommartins@users.noreply.github.com>
Co-authored-by: Luciano Martins <lucianommartins@users.noreply.github.com>
2026-05-27 18:36:27 +08:00
Xin Yang d8eebe6d97 [Perf] Optimize Fp8BlockScaledMMLinearKernel input_scale tensor using new_empty() (#43677)
Signed-off-by: Xin Yang <xyangx@amazon.com>
2026-05-26 19:55:52 -07:00
Jee Jee Li d4004455d2 [Kernel] Remove NormGateLinear (#43554)
Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>
2026-05-25 09:49:19 +00:00
Jee Jee Li b06813e872 [Kernel] Add mhc_pre_big_fuse_with_norm_tilelang (#43474)
Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>
2026-05-25 01:19:45 +00:00
5bb8d2767a [Kernel] Batch invariant NVFP4 linear using cutlass (#39912)
Signed-off-by: Jakub Zakrzewski <jzakrzewski@nvidia.com>
Co-authored-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com>
Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>
2026-05-23 09:41:12 -04:00
Ma JianandKunshang Ji a5bbd81e2e [XPU]feat: enable FP8 block-scaled quantization on XPU (#42952)
Signed-off-by: Ma Jian <jian1.ma@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-05-23 12:33:18 +08:00
Michael Goin 9b54e50e2c [Deprecation] Mark env vars covered by --moe-backend / --linear-backend (#43148)
Signed-off-by: mgoin <mgoin64@gmail.com>
Signed-off-by: Michael Goin <mgoin64@gmail.com>
2026-05-21 12:51:12 -07:00
Wentao YeandIsotr0py 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
Flora Feng 2a43b407c5 [Bugfix][CI] Add missing import of pad_nvfp4_activation_for_cutlass in flashinfer (#43237)
Signed-off-by: sfeng33 <4florafeng@gmail.com>
2026-05-20 11:59:12 -07:00
Wentao YeandMatthew Bonanni 53ff50fcd3 [Perf] Optimize CutlassFP8ScaledMMLinearKernel when padding needed by pre-weight processing, 13.5% TTFT improvement (#42651)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Signed-off-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com>
Co-authored-by: Matthew Bonanni <mbonanni@redhat.com>
2026-05-20 11:57:42 -07:00
meena-at-workandClaude Sonnet 4.6 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
Kunshang Ji 36dcaf25d8 [XPU] add gptq(int4) support (#37844)
Signed-off-by: Kunshang Ji <kunshang.ji@intel.com>
2026-05-19 11:17:09 +08:00
Wentao Ye 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
zofiaandKunshang Ji 0fa888465e [XPU] fix weight scale shape (#42725)
Signed-off-by: Zhu, Zufang <zufang.zhu@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-05-17 16:55:10 +08:00
Michael Goin b2a27b82d9 [Kernel][UX] Add --linear-backend arg for linear kernel selection (#39538)
Signed-off-by: mgoin <mgoin64@gmail.com>
2026-05-15 17:07:39 -07:00
Libin Tangandmergify[bot] <37929162+mergify[bot]@users.noreply.github.com> 9946c38b7f [XPU] Fix double-transpose in XPUFP8ScaledMMLinearKernel for W8A8 quant method (#41689)
Signed-off-by: Libin Tang <libin.tang@intel.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-05-14 17:17:39 +08:00
TJianandYongye Zhu 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
zofiaandKunshang Ji 0ddaf6dffa [XPU] [CT] Enable CT W4A4MxFp4 path and add xpu kernel (#38896)
Signed-off-by: Zhu, Zufang <zufang.zhu@intel.com>
Signed-off-by: zofia <110436990+zufangzhu@users.noreply.github.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-05-13 06:43:00 -07:00
Dipika Sikka a7b801e26d [MXFP4] Support for linear layers + compressed-tensors integration (#41664) 2026-05-12 07:49:33 -04:00
John Calderon 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
Yuwen Zhou 809b98e5b7 [CPU] Add FP8 W8A16 linear support (#41186)
Signed-off-by: yuwenzho <yuwen.zhou@intel.com>
2026-05-06 07:05:27 +00:00
628c436301 [New Model][ROCm] Add AMD support for DeepSeek V4 (#40871)
Signed-off-by: ganyi <ygan@amd.com>
Signed-off-by: whx-sjtu <xiaowang990929@gmail.com>
Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com>
Signed-off-by: tjtanaavllm <tunjian.tan@amd.com>
Co-authored-by: ganyi <ygan@amd.com>
Co-authored-by: tjtanaa <tunjian.tan@embeddedllm.com>
Co-authored-by: tjtanaavllm <tunjian.tan@amd.com>
2026-05-05 08:55:37 -07:00
Andreas Karatzas 5e2c37facd [ROCm][CI] Add missing quantization methods and fix online quant test failures (#39801)
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
2026-04-27 15:08:57 -05:00
+7 4d51588e23 [Feat] DeepSeek V4 Rebased (#40860)
Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai>
Signed-off-by: Woosuk Kwon <woosuk@inferact.ai>
Signed-off-by: qizixi <zixi@inferact.ai>
Signed-off-by: Jee Jee Li <pandaleefree@gmail.com>
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Yongye Zhu <yongye@inferact.ai>
Co-authored-by: Simon Mo <simon@inferact.ai>
Co-authored-by: Bugen Zhao <i@bugenzhao.com>
Co-authored-by: Giancarlo Delfin <gdelfin@inferact.ai>
Co-authored-by: Jee Jee Li <pandaleefree@gmail.com>
Co-authored-by: Nick Hill <nickhill123@gmail.com>
Co-authored-by: Roger Wang <hey@rogerw.io>
Co-authored-by: Roy Wang <yasong.wang@inferact.ai>
Co-authored-by: Woosuk Kwon <woosuk@inferact.ai>
Co-authored-by: youkaichao <youkaichao@gmail.com>
Co-authored-by: Zhewen Li <jerven.vllm@gmail.com>
Co-authored-by: Zijing Liu <liuzijing2014@gmail.com>
Co-authored-by: khluu <khluu000@gmail.com>
Co-authored-by: qizixi <zixi@inferact.ai>
Co-authored-by: Zhewen Li <zhewenli@inferact.ai>
2026-04-26 18:31:08 -07:00
Hashem Hashemi 3a30eaa1d7 Properly enable wvSplitK fp8 path for RDNA (#37712)
Signed-off-by: Hashem Hashemi <hashem.hashemi@amd.com>
2026-04-20 10:09:24 -05:00
nemanjaudovic 58631d7c3f [Bugfix] Fix scaled_mm output narrowing for 3D input tensors (#38093)
Signed-off-by: nemanjaudovic <nudovic@amd.com>
2026-04-20 16:58:39 +08:00
Maral c0c98b8b9a [Bugfix] Add Marlin kernel in block scaled mm kernel selection. (#40105)
Signed-off-by: maral <maralbahari.98@gmail.com>
2026-04-17 10:20:32 +00:00
vllmellm 5f7fab881a [ROCm][FEAT] Integrate aiter gemm w8a8 ptpc (#33773)
Signed-off-by: vllmellm <vllm.ellm@embeddedllm.com>
2026-04-16 09:55:29 +08:00
Wentao Ye 2ad1029233 [Bug] Fix batch invariance nvfp4 support (#39820)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
2026-04-14 17:08:17 -04:00
JartX f414f90601 [Bugfix][Kernel][ROCm] Fix triton_w4a16 scales mismatch when BLOCK_K > group_size (#39705)
Signed-off-by: JartX <sagformas@epdcenter.es>
2026-04-13 14:29:45 -04:00
zofia 1b19bd7589 [MXFP8] [XPU] add a new compressed tensor schema and add a xpu mxfp8 gemm kernel (#38707)
Signed-off-by: Zhu, Zufang <zufang.zhu@intel.com>
2026-04-13 16:59:20 +08:00
Yan Ma 394ff86965 [XPU][CT] support per-channel quantization in xpu fp8 linear method (#38316)
Signed-off-by: Yan Ma <yan.ma@intel.com>
2026-04-12 02:46:28 +00:00
Michael Goin 11e2375fe2 [Refactor] Move MXFP8 GEMM management into MxFp8LinearKernel (#39205)
Signed-off-by: mgoin <mgoin64@gmail.com>
2026-04-10 14:02:03 -07:00
d468322dc1 [Kernel][Hardware][AMD] Add TritonW4A16LinearKernel for ROCm (#37352)
Signed-off-by: jatseng-ai <jatseng@amd.com>
Signed-off-by: jatseng-ai <janet.tseng@amd.com>
Signed-off-by: Matthias Gehre <matthias.gehre@amd.com>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: Matthias Gehre <matthias.gehre@amd.com>
2026-04-10 10:25:27 +00:00