vLLM
/
vllm
mirror of https://github.com/vllm-project/vllm.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity
7,204 Commits 73 Branches 67 Tags 518 MiB
Commit Graph

337 Commits

Author SHA1 Message Date
Sage Moore 558db8083c
[V1][Kernel] Refactor the prefix_prefill kernel so that the caller no longer has to pass in the context lengths (#13095) 2025-02-22 05:25:41 -08:00
Kaixi Hou e109e598c7
[NVIDIA] Support nvfp4 cutlass gemm (#13571) 2025-02-22 05:24:05 -08:00
Lucas Wilkinson 288cc6c234
[Attention] MLA with chunked prefill (#12639) 
Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Patrick Horn <patrick.horn@gmail.com>
Co-authored-by: simon-mo <xmo@berkeley.edu>
Co-authored-by: Tyler Michael Smith <tyler@neuralmagic.com>
 2025-02-21 15:30:12 -08:00
Liangfu Chen 3809458456
[Bugfix] Fix invalid rotary embedding unit test (#13431) 
Signed-off-by: Liangfu Chen <liangfc@amazon.com>
 2025-02-18 11:52:03 +00:00
youkaichao 124776ebd5
[ci] skip failed tests for flashinfer (#13352) 
Signed-off-by: youkaichao <youkaichao@gmail.com>
 2025-02-16 22:09:15 +08:00
wchen61 dc0f7ccf8b
[BugFix] Enhance test_pos_encoding to support execution on multi-devices (#13187) 
Signed-off-by: wchen61 <wchen61@foxmail.com>
 2025-02-16 08:59:49 +00:00
Tyler Michael Smith c1e37bf71b
[Kernel][Bugfix] Refactor and Fix CUTLASS 2:4 Sparse Kernels (#13198) 
Signed-off-by: Tyler Michael Smith <tyler@neuralmagic.com>
 2025-02-14 00:01:14 +00:00
Kaixi Hou 4fc5c23bb6
[NVIDIA] Support nvfp4 quantization (#12784) 2025-02-12 19:51:51 -08:00
Yu Chin Fabian Lim aff404571b
Add Bamba Model (#10909) 
Signed-off-by: Yu Chin Fabian Lim <flim@sg.ibm.com>
Signed-off-by: Tyler Michael Smith <tyler@neuralmagic.com>
Co-authored-by: Tyler Michael Smith <tyler@neuralmagic.com>
 2025-02-06 15:22:42 -08:00
Isotr0py 85ac82d228
[Kernel] Make rotary_embedding ops more flexible with input shape (#12777) 2025-02-06 08:46:13 -08:00
Lucas Wilkinson 75e94309e8
[Perf] Mem align KV caches for CUDA devices (MLA perf improvement) (#12676) 
Signed-off-by: simon-mo <xmo@berkeley.edu>
Signed-off-by: Lucas Wilkinson <lcwilkins@redhat.com>
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
Co-authored-by: simon-mo <xmo@berkeley.edu>
 2025-02-04 18:22:24 -08:00
Hongxia Yang c36ac98d01
[AMD][ROCm] Enable DeepSeek model on ROCm (#12662) 
Signed-off-by: Hongxia Yang <hongxia.yang@amd.com>
Co-authored-by: Matthew Wong <Matthew.Wong2@amd.com>
 2025-02-04 08:24:11 +00:00
Russell Bryant e489ad7a21
[Misc] Add SPDX-License-Identifier headers to python source files (#12628) 
- **Add SPDX license headers to python source files**
- **Check for SPDX headers using pre-commit**

commit 9d7ef44c3cfb72ca4c32e1c677d99259d10d4745
Author: Russell Bryant <rbryant@redhat.com>
Date:   Fri Jan 31 14:18:24 2025 -0500

    Add SPDX license headers to python source files
    
This commit adds SPDX license headers to python source files as
recommended to
the project by the Linux Foundation. These headers provide a concise way
that is
both human and machine readable for communicating license information
for each
source file. It helps avoid any ambiguity about the license of the code
and can
    also be easily used by tools to help manage license compliance.
    
The Linux Foundation runs license scans against the codebase to help
ensure
    we are in compliance with the licenses of the code we use, including
dependencies. Having these headers in place helps that tool do its job.
    
    More information can be found on the SPDX site:
    
    - https://spdx.dev/learn/handling-license-info/
    
    Signed-off-by: Russell Bryant <rbryant@redhat.com>

commit 5a1cf1cb3b80759131c73f6a9dddebccac039dea
Author: Russell Bryant <rbryant@redhat.com>
Date:   Fri Jan 31 14:36:32 2025 -0500

    Check for SPDX headers using pre-commit
    
    Signed-off-by: Russell Bryant <rbryant@redhat.com>

---------

Signed-off-by: Russell Bryant <rbryant@redhat.com>
 2025-02-02 11:58:18 -08:00
Lucas Wilkinson cabaf4eff3
[Attention] MLA decode optimizations (#12528) 
Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
Signed-off-by: simon-mo <xmo@berkeley.edu>
Co-authored-by: Woosuk Kwon <woosuk.kwon@berkeley.edu>
Co-authored-by: simon-mo <simon.mo@hey.com>
Co-authored-by: Michael Goin <mgoin64@gmail.com>
Co-authored-by: Zhuohan Li <zhuohan123@gmail.com>
Co-authored-by: Tyler Michael Smith <tysmith@redhat.com>
Co-authored-by: Alexander Matveev <59768536+alexm-neuralmagic@users.noreply.github.com>
Co-authored-by: simon-mo <xmo@berkeley.edu>
 2025-01-30 23:49:37 -08:00
Lucas Wilkinson 9798b2fb00
[Kernel] Update `cutlass_scaled_mm` to support 2d group (blockwise) scaling (#11868) 2025-01-30 18:33:00 -08:00
Jinzhen Lin 27b78c73ca
[Kernel] add triton fused moe kernel for gptq/awq (#12185) 2025-01-29 09:07:09 -05:00
Harry Mellor 823ab79633
Update `pre-commit` hooks (#12475) 
Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
 2025-01-27 17:23:08 -07:00
Bowen Wang 2bc3fbba0c
[FlashInfer] Upgrade to 0.2.0 (#11194) 
Signed-off-by: Bowen Wang <abmfy@icloud.com>
Signed-off-by: youkaichao <youkaichao@gmail.com>
Co-authored-by: youkaichao <youkaichao@gmail.com>
 2025-01-27 18:19:24 +00:00
Tyler Michael Smith 72f4880425
[Bugfix/CI] Fix broken kernels/test_mha.py (#12450) 2025-01-26 10:39:03 -08:00
Tyler Michael Smith aa2cd2c43d
[Bugfix] Disable w16a16 2of4 sparse CompressedTensors24 (#12417) 
Signed-off-by: Tyler Michael Smith <tyler@neuralmagic.com>
Co-authored-by: mgoin <michael@neuralmagic.com>
 2025-01-26 19:59:58 +08:00
Isotr0py f1fc0510df
[Misc] Add FA2 support to ViT MHA layer (#12355) 
Signed-off-by: Isotr0py <2037008807@qq.com>
 2025-01-25 15:07:35 +08:00
Lucas Wilkinson ab5bbf5ae3
[Bugfix][Kernel] Fix CUDA 11.8 being broken by FA3 build (#12375) 
Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
 2025-01-24 15:27:59 +00:00
Gregory Shtrasberg e97f802b2d
[FP8][Kernel] Dynamic kv cache scaling factors computation (#11906) 
Signed-off-by: Gregory Shtrasberg <Gregory.Shtrasberg@amd.com>
Co-authored-by: Micah Williamson <micah.williamson@amd.com>
 2025-01-23 18:04:03 +00:00
Lucas Wilkinson 978b45f399
[Kernel] Flash Attention 3 Support (#12093) 
Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
 2025-01-23 06:45:48 -08:00
rasmith 68c4421b6d
[AMD][Quantization] Add TritonScaledMMLinearKernel since int8 is broken for AMD (#12282) 
Signed-off-by: Randall Smith <Randall.Smith@amd.com>
 2025-01-23 00:10:37 +00:00
Isotr0py dd7c9ad870
[Bugfix] Remove hardcoded `head_size=256` for Deepseek v2 and v3 (#12067) 
Signed-off-by: Isotr0py <2037008807@qq.com>
 2025-01-16 10:11:54 +00:00
wangxiyuan 3adf0ffda8
[Platform] Do not raise error if _Backend is not found (#12023) 
Signed-off-by: wangxiyuan <wangxiyuan1007@gmail.com>
Signed-off-by: Mengqing Cao <cmq0113@163.com>
Co-authored-by: Mengqing Cao <cmq0113@163.com>
 2025-01-15 10:14:15 +00:00
Jee Jee Li 42f5e7c52a
[Kernel] Support MulAndSilu (#11624) 
Signed-off-by: Jee Jee Li <pandaleefree@gmail.com>
 2025-01-15 02:29:53 +00:00
Avshalom Manevich 263a870ee1
[Hardware][TPU] workaround fix for MoE on TPU (#11764) 2025-01-12 10:53:51 -05:00
Chen Zhang cf5f000d21
[torch.compile] Hide KV cache behind torch.compile boundary (#11677) 
Signed-off-by: Chen Zhang <zhangch99@outlook.com>
 2025-01-10 13:14:42 +08:00
wangxiyuan 405eb8e396
[platform] Allow platform specify attention backend (#11609) 
Signed-off-by: wangxiyuan <wangxiyuan1007@gmail.com>
Signed-off-by: Mengqing Cao <cmq0113@163.com>
Co-authored-by: Mengqing Cao <cmq0113@163.com>
 2025-01-09 21:46:50 +08:00
Chen Zhang e20c92bb61
[Kernel] Move attn_type to Attention.__init__() (#11690) 
Signed-off-by: Chen Zhang <zhangch99@outlook.com>
 2025-01-07 00:11:28 +08:00
Woosuk Kwon 73001445fb
[V1] Implement Cascade Attention (#11635) 
Signed-off-by: Woosuk Kwon <woosuk.kwon@berkeley.edu>
 2025-01-01 21:56:46 +09:00
youkaichao b12e87f942
[platforms] enable platform plugins (#11602) 
Signed-off-by: youkaichao <youkaichao@gmail.com>
 2024-12-30 20:24:45 +08:00
Michael Goin 2072924d14
[Model] [Quantization] Support deepseek_v3 w8a8 fp8 block-wise quantization (#11523) 
Signed-off-by: mgoin <michael@neuralmagic.com>
Signed-off-by: simon-mo <simon.mo@hey.com>
Signed-off-by: simon-mo <xmo@berkeley.edu>
Co-authored-by: simon-mo <simon.mo@hey.com>
Co-authored-by: simon-mo <xmo@berkeley.edu>
Co-authored-by: HandH1998 <1335248067@qq.com>
 2024-12-26 15:33:30 -08:00
Tyler Michael Smith 5a9da2e6e9
[Bugfix][Build/CI] Fix sparse CUTLASS compilation on CUDA [12.0, 12.2) (#11311) 
Signed-off-by: Tyler Michael Smith <tyler@neuralmagic.com>
 2024-12-19 02:43:30 +00:00
Dipika Sikka 60508ffda9
[Kernel]: Cutlass 2:4 Sparsity + FP8/Int8 Quant Support (#10995) 
Co-authored-by: Faraz Shahsavan <faraz.shahsavan@gmail.com>
Co-authored-by: ilmarkov <markovilya197@gmail.com>
Co-authored-by: Rahul Tuli <rahul@neuralmagic.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>
 2024-12-18 09:57:16 -05:00
Luka Govedič 30870b4f66
[torch.compile] Dynamic fp8 + rms_norm fusion (#10906) 
Signed-off-by: luka <luka@neuralmagic.com>
Co-authored-by: Varun Sundar Rabindranath <varun@neuralmagic.com>
 2024-12-13 03:19:23 +00:00
zhou fan 78029b34ed
[BugFix][Kernel]: fix illegal memory access in causal_conv1d when conv_states is None (#10928) 
Signed-off-by: xffxff <1247714429@qq.com>
 2024-12-08 01:21:18 +08:00
Woosuk Kwon 073a4bd1c0
[Kernel] Use `out` arg in flash_attn_varlen_func (#10811) 
Signed-off-by: Woosuk Kwon <woosuk.kwon@berkeley.edu>
 2024-12-01 17:55:39 -08:00
Wallas Henrique c27df94e1f
[Bugfix] Fix chunked prefill with model dtype float32 on Turing Devices (#9850) 
Signed-off-by: Wallas Santos <wallashss@ibm.com>
Co-authored-by: Michael Goin <michael@neuralmagic.com>
 2024-11-25 12:23:32 -05:00
youkaichao 05d1f8c9c6
[misc] move functions to config.py (#10624) 
Signed-off-by: youkaichao <youkaichao@gmail.com>
 2024-11-25 09:27:30 +00:00
youkaichao eebad39f26
[torch.compile] support all attention backends (#10558) 
Signed-off-by: youkaichao <youkaichao@gmail.com>
 2024-11-22 14:04:42 -08:00
Lucas Wilkinson d200972e7f
[Bugfix] Marlin 2:4 temp fix for large M dim (>256) (#10464) 
Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
 2024-11-19 19:40:33 -08:00
ElizaWszola b00b33d77e
[Model][Quantization] HQQ support through Marlin kernel expansion (#9766) 
Signed-off-by: ElizaWszola <eliza@neuralmagic.com>
 2024-11-19 13:31:12 -08:00
Mengqing Cao 8c1fb50705
[Platform][Refactor] Extract func `get_default_attn_backend` to `Platform` (#10358) 
Signed-off-by: Mengqing Cao <cmq0113@163.com>
 2024-11-19 11:22:26 +08:00
Lucas Wilkinson 96d999fbe8
[Kernel] Initial Machete W4A8 support + Refactors (#9855) 
Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
 2024-11-18 12:59:29 -07:00
ElizaWszola 79ee45b428
[Misc] Bump up test_fused_moe tolerance (#10364) 
Signed-off-by: ElizaWszola <eliza@neuralmagic.com>
 2024-11-15 16:31:18 +00:00
Luka Govedič bf2ddc6610
[bugfix] Fix static asymmetric quantization case (#10334) 
Signed-off-by: Daniël de Kok <me@danieldk.eu>
Signed-off-by: luka <luka@neuralmagic.com>
Co-authored-by: Daniël de Kok <me@danieldk.eu>
 2024-11-15 09:35:11 +08:00
rasmith 127c07480e
[Kernel][Triton] Add Triton implementation for scaled_mm_triton to support fp8 and int8 SmoothQuant, symmetric case (#9857) 
Signed-off-by: Randall Smith <Randall.Smith@amd.com>
 2024-11-08 19:59:22 -05:00
Luka Govedič 4f93dfe952
[torch.compile] Fuse RMSNorm with quant (#9138) 
Signed-off-by: luka <luka@neuralmagic.com>
Co-authored-by: youkaichao <youkaichao@126.com>
 2024-11-08 21:20:08 +00:00
Joe Runde d58268c56a
[V1] Make v1 more testable (#9888) 
Signed-off-by: Joe Runde <Joseph.Runde@ibm.com>
 2024-11-06 11:57:35 -08:00
Aaron Pham 21063c11c7
[CI/Build] drop support for Python 3.8 EOL (#8464) 
Signed-off-by: Aaron Pham <contact@aarnphm.xyz>
 2024-11-06 07:11:55 +00:00
Michael Goin 235366fe2e
[CI] Prune back the number of tests in tests/kernels/* (#9932) 
Signed-off-by: mgoin <michael@neuralmagic.com>
 2024-11-05 16:02:32 -05:00
sroy745 a78dd3303e
[Encoder Decoder] Add flash_attn kernel support for encoder-decoder models (#9559) 2024-11-01 23:22:49 -07:00
Peter Salas 6c0b7f548d
[Core][VLM] Add precise multi-modal placeholder tracking (#8346) 
Signed-off-by: Peter Salas <peter@fixie.ai>
 2024-11-01 16:21:10 -07:00
Pavani Majety 598b6d7b07
[Bugfix/Core] Flashinfer k_scale and v_scale (#9861) 2024-11-01 12:15:05 -07:00
Mor Zusman 9fb12f7848
[BugFix][Kernel] Fix Illegal memory access in causal_conv1d in H100 (#9838) 
Signed-off-by: mzusman <mor.zusmann@gmail.com>
 2024-10-31 20:06:25 +00:00
wangshuai09 622b7ab955
[Hardware] using current_platform.seed_everything (#9785) 
Signed-off-by: wangshuai09 <391746016@qq.com>
 2024-10-29 14:47:44 +00:00
youkaichao 32176fee73
[torch.compile] support moe models (#9632) 
Signed-off-by: youkaichao <youkaichao@gmail.com>
 2024-10-27 21:58:04 -07:00
wangshuai09 4e2d95e372
[Hardware][ROCM] using current_platform.is_rocm (#9642) 
Signed-off-by: wangshuai09 <391746016@qq.com>
 2024-10-28 04:07:00 +00:00
Mengqing Cao 5cbdccd151
[Hardware][openvino] is_openvino --> current_platform.is_openvino (#9716) 2024-10-26 10:59:06 +00:00
Charlie Fu 59449095ab
[Performance][Kernel] Fused_moe Performance Improvement (#9384) 
Signed-off-by: charlifu <charlifu@amd.com>
 2024-10-24 15:37:52 -07:00
Jee Jee Li 295a061fb3
[Kernel] add kernel for FATReLU (#9610) 
Signed-off-by: Jee Jee Li <pandaleefree@gmail.com>
 2024-10-24 16:18:27 +08:00
wangshuai09 3ddbe25502
[Hardware][CPU] using current_platform.is_cpu (#9536) 2024-10-22 00:50:43 -07:00
Chen Zhang 4fa3e33349
[Kernel] Support sliding window in flash attention backend (#9403) 2024-10-20 10:57:52 -07:00
bnellnm eca2c5f7c0
[Bugfix] Fix support for dimension like integers and ScalarType (#9299) 2024-10-17 19:08:34 +00:00
Mor Zusman fb60ae9b91
[Kernel][Model] Improve continuous batching for Jamba and Mamba (#9189) 2024-10-16 12:12:43 -04:00
Cyrus Leung 7e7eae338d
[Misc] Standardize RoPE handling for Qwen2-VL (#9250) 2024-10-16 13:56:17 +08:00
Tyler Michael Smith 7342a7d7f8
[Model] Support Mamba (#6484) 2024-10-11 15:40:06 +00:00
Lucas Wilkinson a64e7b9407
[Bugfix] Machete garbage results for some models (large K dim) (#9212) 2024-10-10 14:16:17 +08:00
bnellnm bd37b9fbe2
[Bugfix] Try to handle older versions of pytorch (#9086) 2024-10-08 14:28:12 -07:00
ElizaWszola 05d686432f
[Kernel] Zero point support in fused MarlinMoE kernel + AWQ Fused MoE (#8973) 
Co-authored-by: Dipika <dipikasikka1@gmail.com>
Co-authored-by: Dipika Sikka <ds3822@columbia.edu>
 2024-10-04 12:34:44 -06:00
youkaichao 9aaf14c62e
[misc] add forward context for attention (#9029) 2024-10-03 12:09:42 -07:00
Mor Zusman f13a07b1f8
[Kernel][Model] Varlen prefill + Prefill chunking support for mamba kernels and Jamba model (#8533) 2024-09-29 17:35:58 -04:00
ElizaWszola d081da0064
[Bugfix] Fix Marlin MoE act order when is_k_full == False (#8741) 
Co-authored-by: Tyler Michael Smith <tyler@neuralmagic.com>
 2024-09-28 18:19:40 -07:00
youkaichao a9b15c606f
[torch.compile] use empty tensor instead of None for profiling (#8875) 2024-09-27 08:11:32 -07:00
bnellnm 300da09177
[Kernel] Fullgraph and opcheck tests (#8479) 2024-09-25 08:35:52 -06:00
Lucas Wilkinson 86e9c8df29
[Kernel] (2/N) Machete - Integrate into CompressedTensorsWNA16 and GPTQMarlin (#7701) 
Co-authored-by: mgoin <michael@neuralmagic.com>
Co-authored-by: Divakar Verma <137818590+divakar-amd@users.noreply.github.com>
Co-authored-by: Tyler Michael Smith <tyler@neuralmagic.com>
 2024-09-23 13:46:26 -04:00
Charlie Fu 9cc373f390
[Kernel][Amd] Add fp8 kv cache support for rocm custom paged attention (#8577) 2024-09-19 17:37:57 +00:00
Tyler Michael Smith db9120cded
[Kernel] Change interface to Mamba selective_state_update for continuous batching (#8039) 2024-09-18 20:05:06 +00:00
Cyrus Leung 6ffa3f314c
[CI/Build] Avoid CUDA initialization (#8534) 2024-09-18 10:38:11 +00:00
Tyler Michael Smith 8110e44529
[Kernel] Change interface to Mamba causal_conv1d_update for continuous batching (#8012) 2024-09-17 23:44:27 +00:00
Simon Mo 546034b466
[refactor] remove triton based sampler (#8524) 2024-09-16 20:04:48 -07:00
Luka Govedič 5d73ae49d6
[Kernel] AQ AZP 3/4: Asymmetric quantization kernels (#7270) 2024-09-16 11:52:40 -07:00
ElizaWszola a091e2da3e
[Kernel] Enable 8-bit weights in Fused Marlin MoE (#8032) 
Co-authored-by: Dipika <dipikasikka1@gmail.com>
 2024-09-16 09:47:19 -06:00
Isotr0py fc990f9795
[Bugfix][Kernel] Add `IQ1_M` quantization implementation to GGUF kernel (#8357) 2024-09-15 16:51:44 -06:00
Charlie Fu 1ef0d2efd0
[Kernel][Hardware][Amd]Custom paged attention kernel for rocm (#8310) 2024-09-13 17:01:11 -07:00
Cyrus Leung a84e598e21
[CI/Build] Reorganize models tests (#7820) 2024-09-13 10:20:06 -07:00
bnellnm 73202dbe77
[Kernel][Misc] register ops to prevent graph breaks (#6917) 
Co-authored-by: Sage Moore <sage@neuralmagic.com>
 2024-09-11 12:52:19 -07:00
Dipika Sikka 6cd5e5b07e
[Misc] Fused MoE Marlin support for GPTQ (#8217) 2024-09-09 23:02:52 -04:00
Elfie Guo e39ebf5cf5
[Core/Bugfix] Add query dtype as per FlashInfer API requirements. (#8173) 2024-09-05 05:12:26 +00:00
Pavani Majety 6b3421567d
[Core][Kernels] Enable FP8 KV Cache with Flashinfer backend. + BugFix for kv_cache_dtype=auto (#7985) 
Co-authored-by: Simon Mo <simon.mo@hey.com>
Co-authored-by: Cody Yu <hao.yu.cody@gmail.com>
 2024-08-29 14:53:11 -04:00
youkaichao ef99a78760
Revert "[Core][Kernels] Use FlashInfer backend for FP8 KV Cache when available." (#7982) 2024-08-28 21:27:06 -07:00
Mor Zusman fdd9daafa3
[Kernel/Model] Migrate mamba_ssm and causal_conv1d kernels to vLLM (#7651) 2024-08-28 15:06:52 -07:00
rasmith e5697d161c
[Kernel] [Triton] [AMD] Adding Triton implementations awq_dequantize and awq_gemm to support AWQ (#7386) 2024-08-28 15:37:47 -04:00
Pavani Majety b98cc28f91
[Core][Kernels] Use FlashInfer backend for FP8 KV Cache when available. (#7798) 
Co-authored-by: Simon Mo <simon.mo@hey.com>
 2024-08-28 10:01:22 -07:00
LI MOU 53328d7536
[BUG] fix crash on flashinfer backend with cudagraph disabled, when attention group_size not in [1,2,4,8] (#7509) 2024-08-21 08:54:31 -07:00
Lucas Wilkinson 5288c06aa0
[Kernel] (1/N) Machete - Hopper Optimized Mixed Precision Linear Kernel (#7174) 2024-08-20 07:09:33 -06:00
Charlie Fu e837b624f2
[Feature][Hardware][Amd] Add fp8 Linear Layer for Rocm (#7210) 2024-08-16 10:06:30 -07:00
youkaichao 54bd9a03c4
register custom op for flash attn and use from torch.ops (#7536) 2024-08-15 22:38:56 -07:00
jon-chuang 50b8d08dbd
[Misc/Testing] Use `torch.testing.assert_close` (#7324) 2024-08-16 04:24:04 +00:00
jon-chuang a046f86397
[Core/Bugfix] Add FP8 K/V Scale and dtype conversion for prefix/prefill Triton Kernel (#7208) 
Co-authored-by: Cody Yu <hao.yu.cody@gmail.com>
 2024-08-12 22:47:41 +00:00
Luka Govedič 5fb4a3f678
[Bugfix][Kernel] Increased atol to fix failing tests (#7305) 2024-08-08 12:16:13 -04:00
afeldman-nm fd95e026e0
[Core] Subclass ModelRunner to support cross-attention & encoder sequences (towards eventual encoder/decoder model support) (#4942) 
Co-authored-by: Andrew Feldman <afeld2012@gmail.com>
Co-authored-by: Nick Hill <nickhill@us.ibm.com>
 2024-08-06 16:51:47 -04:00
Luka Govedič 8d59dbb000
[Kernel] Add per-tensor and per-token AZP epilogues (#5941) 
Co-authored-by: Tyler Michael Smith <tyler@neuralmagic.com>
 2024-08-06 18:17:08 +00:00
Lucas Wilkinson a8d604ca2a
[Misc] Disambiguate quantized types via a new ScalarType (#6396) 2024-08-02 13:51:58 -07:00
Woosuk Kwon 805a8a75f2
[Misc] Support attention logits soft-capping with flash-attn (#7022) 2024-08-01 13:14:37 -07:00
Jee Jee Li 7ecee34321
[Kernel][RFC] Refactor the punica kernel based on Triton (#5036) 2024-07-31 17:12:24 -07:00
HandH1998 6512937de1
Support W4A8 quantization for vllm (#5218) 2024-07-31 07:55:21 -06:00
Varun Sundar Rabindranath af647fb8b3
[Kernel] Tuned int8 kernels for Ada Lovelace (#6848) 
Co-authored-by: Varun Sundar Rabindranath <varun@neuralmagic.com>
 2024-07-29 20:24:58 -06:00
Thomas Parnell 9a7e2d0534
[Bugfix] Allow vllm to still work if triton is not installed. (#6786) 
Signed-off-by: Thomas Parnell <tpa@zurich.ibm.com>
 2024-07-29 14:51:27 -07:00
Varun Sundar Rabindranath 766435e660
[Kernel] Tuned FP8 Kernels for Ada Lovelace (#6677) 
Co-authored-by: Varun Sundar Rabindranath <varun@neuralmagic.com>
 2024-07-29 09:42:35 -06:00
Alexander Matveev 75acdaa4b6
[Kernel] Increase precision of GPTQ/AWQ Marlin kernel (#6795) 2024-07-27 17:52:33 -04:00
Joe 14dbd5a767
[Model] H2O Danube3-4b (#6451) 2024-07-26 20:47:50 -07:00
Antoni Baum 0e63494cf3
Add fp8 support to `reshape_and_cache_flash` (#6667) 2024-07-24 18:36:52 +00:00
Tyler Michael Smith fea59c7712
[Bugfix][Kernel] Use int64_t for indices in fp8 quant kernels (#6649) 2024-07-22 14:08:30 -06:00
Alexander Matveev 396d92d5e0
[Kernel][Core] Add AWQ support to the Marlin kernel (#6612) 2024-07-21 19:41:42 -04:00
Varun Sundar Rabindranath 2e26564259
[ Kernel ] FP8 Dynamic Per Token Quant - Add scale_ub (#6593) 
Co-authored-by: Varun Sundar Rabindranth <varun@neuralmagic.com>
 2024-07-19 18:15:26 -07:00
Robert Shaw 4cc24f01b1
[ Kernel ] Enable Dynamic Per Token `fp8` (#6547) 2024-07-19 23:08:15 +00:00
Varun Sundar Rabindranath b5241e41d9
[ Kernel ] FP8 Dynamic-Per-Token Quant Kernel (#6511) 
Co-authored-by: Varun Sundar Rabindranath <varun@neuralmagic.com>
 2024-07-18 01:38:35 +00:00
Michael Goin 978aed5300
[Kernel][Attention] Separate `Attention.kv_scale` into `k_scale` and `v_scale` (#6081) 2024-07-16 15:31:32 -07:00
Robert Shaw b675069d74
[ Misc ] Refactor Marlin Python Utilities (#6082) 
Co-authored-by: Robert Shaw <rshaw@neuralmagic.com>
 2024-07-11 15:40:11 +00:00
afeldman-nm 543aa48573
[Kernel] Correctly invoke prefill & decode kernels for cross-attention (towards eventual encoder/decoder model support) (#4888) 
Co-authored-by: Woosuk Kwon <woosuk.kwon@berkeley.edu>
 2024-07-08 17:12:15 +00:00
Lily Liu 69ec3ca14c
[Kernel][Model] logits_soft_cap for Gemma2 with flashinfer (#6051) 
Co-authored-by: Simon Mo <simon.mo@hey.com>
 2024-07-04 16:35:51 -07:00
Michael Goin 47f0954af0
[Kernel] Expand FP8 support to Ampere GPUs using FP8 Marlin (#5975) 2024-07-03 17:38:00 +00:00
youkaichao 482045ee77
[hardware][misc] introduce platform abstraction (#6080) 2024-07-02 20:12:22 -07:00
Robert Shaw 7c008c51a9
[ Misc ] Refactor MoE to isolate Fp8 From Mixtral (#5970) 
Co-authored-by: Robert Shaw <rshaw@neuralmagic>
Co-authored-by: Michael Goin <michael@neuralmagic.com>
 2024-07-02 21:54:35 +00:00
Avshalom Manevich 12a59959ed
[Bugfix] adding chunking mechanism to fused_moe to handle large inputs (#6029) 2024-07-01 21:08:29 +00:00
youkaichao 614aa51203
[misc][cuda] use nvml to avoid accidentally cuda initialization (#6007) 2024-06-30 20:07:34 -07:00
Tyler Michael Smith 6a2d659d28
[Bugfix] Fix compute datatype for cutlass 3.x epilogues (#5931) 2024-06-28 17:10:34 +00:00
Ilya Lavrenov 57f09a419c
[Hardware][Intel] OpenVINO vLLM backend (#5379) 2024-06-28 13:50:16 +00:00
Luka Govedič 5bfd1bbc98
[Kernel] Adding bias epilogue support for `cutlass_scaled_mm` (#5560) 
Co-authored-by: Chih-Chieh-Yang <7364402+cyang49@users.noreply.github.com>
Co-authored-by: Lucas Wilkinson <lwilkinson@neuralmagic.com>
 2024-06-26 15:16:00 +00:00
Cyrus Leung 0e9164b40a
[mypy] Enable type checking for test directory (#5017) 2024-06-15 04:45:31 +00:00
Tyler Michael Smith 85657b5607
[Kernel] Factor out epilogues from cutlass kernels (#5391) 
Co-authored-by: Michael Goin <michael@neuralmagic.com>
Co-authored-by: youkaichao <youkaichao@gmail.com>
Co-authored-by: zifeitong <zifei.tong@parasail.io>
Co-authored-by: Robert Shaw <114415538+robertgshaw2-neuralmagic@users.noreply.github.com>
 2024-06-13 11:22:19 -07:00
bnellnm 5467ac3196
[Kernel][Misc] Use TORCH_LIBRARY instead of PYBIND11_MODULE for custom ops (#5047) 2024-06-09 16:23:30 -04:00
Dipika Sikka ca3ea51bde
[Kernel] Dynamic Per-Token Activation Quantization (#5037) 
Co-authored-by: Varun Sundar Rabindranath <varunsundar08@gmail.com>
Co-authored-by: Varun Sundar Rabindranath <varun@neuralmagic.com>
 2024-06-07 09:36:26 -07:00
Woosuk Kwon 41ca62cf03
[Misc] Add CustomOp interface for device portability (#5255) 2024-06-05 09:18:19 -07:00
afeldman-nm f42a006b15
[Bugfix]: During testing, use pytest monkeypatch for safely overriding the env var that indicates the vLLM backend (#5210) 2024-06-03 20:32:57 -07:00
Tyler Michael Smith cbb2f59cc8
[Kernel] Pass a device pointer into the quantize kernel for the scales (#5159) 2024-06-03 09:52:30 -07:00
Varun Sundar Rabindranath f081c3ce4b
[Kernel] Update Cutlass fp8 configs (#5144) 
Co-authored-by: Varun Sundar Rabindranath <varun@neuralmagic.com>
Co-authored-by: Robert Shaw <114415538+robertgshaw2-neuralmagic@users.noreply.github.com>
 2024-06-01 08:46:07 +00:00
Tyler Michael Smith 260d119e86
[Kernel] Refactor CUTLASS kernels to always take scales that reside on the GPU (#5137) 2024-06-01 06:45:32 +00:00
SnowDist a22dea54d3
[Model] Support MAP-NEO model (#5081) 
Co-authored-by: Zhuohan Li <zhuohan123@gmail.com>
 2024-05-30 19:24:41 -07:00
Eric Xihui Lin 8e192ff967
[Kernel][Backend][Model] Blocksparse flash attention kernel and Phi-3-Small model (#4799) 
Co-authored-by: beagleski <yunanzhang@microsoft.com>
Co-authored-by: bapatra <bapatra@microsoft.com>
Co-authored-by: Barun Patra <codedecde@users.noreply.github.com>
Co-authored-by: Michael Goin <michael@neuralmagic.com>
 2024-05-24 22:00:52 -07:00
Dipika Sikka a1242324c9
[Kernel] Initial Activation Quantization Support (#4525) 
Co-authored-by: Varun Sundar Rabindranath <varunsundar08@gmail.com>
Co-authored-by: Varun Sundar Rabindranath <varun@neuralmagic.com>
 2024-05-23 21:29:18 +00:00
Alexander Matveev 6066253296
Marlin 24 prefill performance improvement (about 25% better on average) (#4983) 2024-05-23 02:39:27 -04:00
Cody Yu ee3eea0a1b
[Misc] Take user preference in attention selector (#4960) 2024-05-23 07:55:56 +09:00
Tyler Michael Smith 8674f9880e
[Kernel] Fixup for CUTLASS kernels in CUDA graphs (#4954) 
Pass the CUDA stream into the CUTLASS GEMMs, to avoid future issues with CUDA graphs
 2024-05-22 14:10:43 +00:00
Woosuk Kwon b57e6c5949
[Kernel] Add flash-attn back (#4907) 2024-05-19 18:11:30 -07:00
Alexander Matveev 27ce85476e
[Kernel] Add marlin_24 unit tests (#4901) 2024-05-19 11:37:34 -04:00