iHore/vLLM
8
0
Fork 0
Code Issues Pull Requests Actions 2 Packages Projects Releases Wiki Activity
1,863 Commits 3 Branches 0 Tags
fc17110bbef4e78703abffac51133a2fb71e9f79
Commit Graph
73 Commits
Author SHA1 Message Date
Robert ShawandRobert Shaw abfe705a02 [ Misc ] Support Fp8 via llm-compressor (#6110) 
Co-authored-by: Robert Shaw <rshaw@neuralmagic>
 2024-07-07 20:42:11 +00:00
Robert Shaw 62963d129e [ Misc ] Clean Up CompressedTensorsW8A8 (#6113) 2024-07-03 22:50:08 +00: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
Qubitium-ModelCloudRobert ShawZX
 		ee93f4f92a [CORE] Quantized lm-head Framework (#4442) 
Co-authored-by: Robert Shaw <rshaw@neuralmagic.com>
Co-authored-by: ZX <zx@lbx.dev>
 2024-07-02 22:25:17 +00:00
Robert ShawRobert ShawMichael Goin
 		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
youkaichao 614aa51203 [misc][cuda] use nvml to avoid accidentally cuda initialization (#6007) 2024-06-30 20:07:34 -07:00
Robert ShawandRobert Shaw af9ad46fca [ Misc ] Refactor w8a8 to use process_weights_after_load (Simplify Weight Loading) (#5940) 
Co-authored-by: Robert Shaw <rshaw@neuralmagic>
 2024-06-30 23:06:27 +00:00
Dipika Sikka 7836fdcc11 [Misc] Fix get_min_capability (#5971) 2024-06-30 20:15:16 +00:00
Robert ShawMichael GoinRobert Shaw
 		8dbfcd35bf [ CI/Build ] Added E2E Test For Compressed Tensors (#5839) 
Co-authored-by: Michael Goin <michael@neuralmagic.com>
Co-authored-by: Robert Shaw <rshaw@neuralmagic>
 2024-06-29 21:12:58 +08:00
Robert ShawandRobert Shaw 2cd402e169 [ Bugfix ] Enabling Loading Models With Fused QKV/MLP on Disk with FP8 (#5921) 
Co-authored-by: Robert Shaw <rshaw@neuralmagic>
 2024-06-28 18:43:49 +00:00
Dipika Sikka dd248f7675 [Misc] Update w4a16 compressed-tensors support to include w8a16 (#5794) 2024-06-25 19:23:35 +00:00
Tyler Michael Smith 3f3b6b2150 [Bugfix] Fix the CUDA version check for FP8 support in the CUTLASS kernels (#5715) 2024-06-20 18:36:10 +00:00
Dipika Sikka 4a30d7e3cc [Misc] Add per channel support for static activation quantization; update w8a8 schemes to share base classes (#5650) 2024-06-19 18:06:44 -04:00
Dipika Sikka 95db455e7f [Misc] Add channel-wise quantization support for w8a8 dynamic per token activation quantization (#5542) 2024-06-18 12:45:05 -04:00
Dipika Sikka 890d8d960b [Kernel] compressed-tensors marlin 24 support (#5435) 2024-06-17 12:32:48 -04:00
Cyrus Leung 0e9164b40a [mypy] Enable type checking for test directory (#5017) 2024-06-15 04:45:31 +00:00
Robert ShawmgoinDipika Sikka
 		15985680e2 [ Misc ] Rs/compressed tensors cleanup (#5432) 
Co-authored-by: mgoin <michael@neuralmagic.com>
Co-authored-by: Dipika Sikka <dipikasikka1@gmail.com>
 2024-06-14 10:01:46 -07:00
Tyler Michael Smith 703475f6c2 [Kernel] Fix CUTLASS 3.x custom broadcast load epilogue (#5516) 2024-06-14 09:30:15 -07:00
Tyler Michael Smith e38042d4af [Kernel] Disable CUTLASS kernels for fp8 (#5505) 2024-06-13 13:38:05 -07:00
Tyler Michael SmithMichael GoinyoukaichaozifeitongRobert Shaw
 		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
Dipika SikkaandRobert Shaw c2637a613b [Kernel] w4a16 support for compressed-tensors (#5385) 
Co-authored-by: Robert Shaw <114415538+robertgshaw2-neuralmagic@users.noreply.github.com>
 2024-06-13 10:19:56 -04:00
Dipika SikkaandMichael Goin 5884c2b454 [Misc] Update to comply with the new compressed-tensors config (#5350) 
Co-authored-by: Michael Goin <michael@neuralmagic.com>
 2024-06-10 03:49:46 +00:00
Michael Goin c09dade2a2 [Misc][Breaking] Change FP8 checkpoint format from act_scale -> input_scale (#5353) 2024-06-08 13:54:05 -04:00
Cheng Li e69ded7d1c [Bug Fix] Fix the support check for FP8 CUTLASS (#5352) 
Bug description:
With torch 2.4.0.dev20240603+cu121,
cutlass_fp8_supported outputs False, and the (capability, version) before the comparison is (90, 11111111112)

This PR fixes the support check for FP8 CUTLASS ( cutlass_fp8_supported) which was introduced in https://github.com/vllm-project/vllm/pull/5183.
 2024-06-08 00:42:05 +00:00
Dipika SikkaVarun Sundar RabindranathVarun Sundar Rabindranath
 		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
Tyler Michael Smith 8d75fe48ca [Kernel] Switch fp8 layers to use the CUTLASS kernels (#5183) 
Switching from torch._scaled_mm to vLLM's cutlass fp8 kernels when supported as we are seeing 5-15% improvement in e2e performance on neuralmagic/Meta-Llama-3-8B-Instruct-FP8

see https://docs.google.com/spreadsheets/d/1GiAnmzyGHgZ6zL_LDSTm35Bdrt4A8AaFEurDlISYYA4/ for some quick e2e benchmarks and #5144 for comparisons across different GEMM sizes.
 2024-06-07 08:42:35 +00:00
Cody Yu 5563a4dea8 [Model] Correct Mixtral FP8 checkpoint loading (#5231) 2024-06-05 10:58:50 -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
chenqianfzh b9c0605a8e [Feature][Kernel] Support bitsandbytes quantization and QLoRA (#4776) 2024-06-01 14:51:10 -06: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
Robert Shaw b35be5403f [Bugfix] Avoid Warnings in SparseML Activation Quantization (#5120) 2024-05-30 17:04:37 -07:00
Alexander Matveev 5bf185a1c4 [Bugfix] gptq_marlin: Ensure g_idx_sort_indices is not a Parameter (#5108) 2024-05-30 00:30:18 +00:00
sasha0552 fbdb7b3ee2 [Core] Allow AQLM on Pascal (#5058) 2024-05-27 15:26:14 -07:00
Dipika SikkaVarun Sundar RabindranathVarun Sundar Rabindranath
 		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 a3a73ab069 [Misc] Load FP8 kv-cache scaling factors from checkpoints (#4893) 
The 2nd PR for #4532.

This PR supports loading FP8 kv-cache scaling factors from a FP8 checkpoint (with .kv_scale parameter).
 2024-05-22 13:28:20 -07:00
Alexander Matveev 27ce85476e [Kernel] Add marlin_24 unit tests (#4901) 2024-05-19 11:37:34 -04:00
Alexander MatveevandRobert Shaw 6979ade384 Add GPTQ Marlin 2:4 sparse structured support (#4790) 
Co-authored-by: Robert Shaw <rshaw@neuralmagic.com>
 2024-05-16 12:56:15 -04:00
Jinzhen Lin 99caa49106 [Kernel] add bfloat16 support for gptq marlin kernel (#4788) 2024-05-16 09:55:29 -04:00
alexm-nm 5c342570d7 Add marlin unit tests and marlin benchmark script (#4815) 2024-05-16 09:36:49 -04:00
Philipp Moritz 379da6dcb5 [Kernel] [FP8] Improve FP8 linear layer performance (#4691) 
This PR improves the FP8 performance of linear layers, which had been lacking before (#4118 (comment) and #4118 (comment)).

We noticed that CUBLASLt can find a better algorithm if the first dimension of the matrix is greater than 16. So this PR enlarges matrices appropriately during quantization. This improves FP8 performance and removes the performance regression vs. FP16, in many cases exceeding FP16 performance.

Here are benchmarks on llama3 70b (ITL numbers for 1000 input and 50 output tokens at fixed qps and at TP 4), all FP8 measurements are for dynamic quantization:

qps = 1: 24 ms (FP8, this PR), 32 ms (FP8, previous main), 26 ms (FP16)
qps = 2: 26 ms (FP8, this PR), 34ms (FP8, previous main), 28 ms (FP16) 
qps = 4: 33 ms (FP8, this PR), 44 ms (FP8, previous main), 36 ms (FP16)
qps = 6: 46 ms (FP8, this PR), 56 ms (FP8, previous main), 54 ms (FP16)
qps = 8: 85 ms (FP8, this PR), 85 ms (FP8, previous main), 138 ms (FP16)
 2024-05-09 16:38:07 -07:00
Hao ZhangDash DesaiAurick QiaoAurick QiaoAurick QiaoCody Yu
 		ebce310b74 [Model] Snowflake arctic model implementation (#4652) 
Co-authored-by: Dash Desai <1723932+iamontheinet@users.noreply.github.com>
Co-authored-by: Aurick Qiao <qiao@aurick.net>
Co-authored-by: Aurick Qiao <aurick.qiao@snowflake.com>
Co-authored-by: Aurick Qiao <aurickq@users.noreply.github.com>
Co-authored-by: Cody Yu <hao.yu.cody@gmail.com>
 2024-05-09 22:37:14 +00:00
alexm-nm 7038e8b803 [Kernel] Support running GPTQ 8-bit models in Marlin (#4533) 2024-05-02 12:56:22 -04:00
Robert ShawPhilipp MoritzWoosuk KwonmgoinTyler Michael SmithCody Yu
 		111815d482 [Kernel] Support Fp8 Checkpoints (Dynamic + Static) (#4332) 
Co-authored-by: Philipp Moritz <pcmoritz@gmail.com>
Co-authored-by: Woosuk Kwon <woosuk.kwon@berkeley.edu>
Co-authored-by: mgoin <michael@neuralmagic.com>
Co-authored-by: Tyler Michael Smith <tyler@neuralmagic.com>
Co-authored-by: Cody Yu <hao.yu.cody@gmail.com>
 2024-04-30 21:46:12 +00:00
Kunshang Ji 26f2fb5113 [Core]Refactor gptq_marlin ops (#4466) 2024-04-30 08:14:47 -04:00
Robert Shawalexmmgoin
 		73c8d677e5 [Kernel] Marlin Expansion: Support AutoGPTQ Models with Marlin (#3922) 
Co-authored-by: alexm <alexm@neuralmagic.com>
Co-authored-by: mgoin <michael@neuralmagic.com>
 2024-04-29 09:35:34 -07:00
SangBin Cho df29793dc7 [mypy][5/N] Support all typing on model executor (#4427) 2024-04-28 19:01:26 -07:00
Philipp MoritzandWoosuk Kwon 12628d3c78 [Kernel] Optimize FP8 support for MoE kernel / Mixtral via static scales (#4343) 
Co-authored-by: Woosuk Kwon <woosuk.kwon@berkeley.edu>
 2024-04-27 04:49:59 +00:00
Cody Yu a62aaf1df5 [Misc][Refactor] Generalize linear_method to be quant_method (#4373) 2024-04-26 16:41:14 -04:00