Previous News
+- [2025/08] We hosted [vLLM Korea Meetup](https://luma.com/cgcgprmh) with Red Hat and Rebellions! We shared the latest advancements in vLLM along with project spotlights from the vLLM Korea community. Please find the meetup slides [here](https://drive.google.com/file/d/1bcrrAE1rxUgx0mjIeOWT6hNe2RefC5Hm/view).
+- [2025/08] We hosted [vLLM Beijing Meetup](https://mp.weixin.qq.com/s/dgkWg1WFpWGO2jCdTqQHxA) focusing on large-scale LLM deployment! Please find the meetup slides [here](https://drive.google.com/drive/folders/1Pid6NSFLU43DZRi0EaTcPgXsAzDvbBqF) and the recording [here](https://www.chaspark.com/#/live/1166916873711665152).
+- [2025/05] We hosted [NYC vLLM Meetup](https://lu.ma/c1rqyf1f)! Please find the meetup slides [here](https://docs.google.com/presentation/d/1_q_aW_ioMJWUImf1s1YM-ZhjXz8cUeL0IJvaquOYBeA/edit?usp=sharing).
+- [2025/04] We hosted [Asia Developer Day](https://www.sginnovate.com/event/limited-availability-morning-evening-slots-remaining-inaugural-vllm-asia-developer-day)! Please find the meetup slides from the vLLM team [here](https://docs.google.com/presentation/d/19cp6Qu8u48ihB91A064XfaXruNYiBOUKrBxAmDOllOo/edit?usp=sharing).
- [2025/03] We hosted [vLLM x Ollama Inference Night](https://lu.ma/vllm-ollama)! Please find the meetup slides from the vLLM team [here](https://docs.google.com/presentation/d/16T2PDD1YwRnZ4Tu8Q5r6n53c5Lr5c73UV9Vd2_eBo4U/edit?usp=sharing).
- [2025/03] We hosted [the first vLLM China Meetup](https://mp.weixin.qq.com/s/n77GibL2corAtQHtVEAzfg)! Please find the meetup slides from vLLM team [here](https://docs.google.com/presentation/d/1REHvfQMKGnvz6p3Fd23HhSO4c8j5WPGZV0bKYLwnHyQ/edit?usp=sharing).
- [2025/03] We hosted [the East Coast vLLM Meetup](https://lu.ma/7mu4k4xx)! Please find the meetup slides [here](https://docs.google.com/presentation/d/1NHiv8EUFF1NLd3fEYODm56nDmL26lEeXCaDgyDlTsRs/edit#slide=id.g31441846c39_0_0).
@@ -121,6 +126,7 @@ Cash Donations:
Compute Resources:
+- Alibaba Cloud
- AMD
- Anyscale
- AWS
@@ -160,7 +166,7 @@ If you use vLLM for your research, please cite our [paper](https://arxiv.org/abs
## Contact Us
-- For technical questions and feature requests, please use GitHub [Issues](https://github.com/vllm-project/vllm/issues) or [Discussions](https://github.com/vllm-project/vllm/discussions)
+- For technical questions and feature requests, please use GitHub [Issues](https://github.com/vllm-project/vllm/issues)
- For discussing with fellow users, please use the [vLLM Forum](https://discuss.vllm.ai)
- For coordinating contributions and development, please use [Slack](https://slack.vllm.ai)
- For security disclosures, please use GitHub's [Security Advisories](https://github.com/vllm-project/vllm/security/advisories) feature
diff --git a/SECURITY.md b/SECURITY.md
index 414669fb37..d6319cdb1a 100644
--- a/SECURITY.md
+++ b/SECURITY.md
@@ -42,4 +42,9 @@ For certain security issues of CRITICAL, HIGH, or MODERATE severity level, we ma
* If you wish to be added to the prenotification group, please send an email copying all the members of the [vulnerability management team](https://docs.vllm.ai/en/latest/contributing/vulnerability_management.html). Each vendor contact will be analyzed on a case-by-case basis.
+* Organizations and vendors who either ship or use vLLM, are eligible to join the prenotification group if they meet at least one of the following qualifications
+ * Substantial internal deployment leveraging the upstream vLLM project.
+ * Established internal security teams and comprehensive compliance measures.
+ * Active and consistent contributions to the upstream vLLM project.
+
* We may withdraw organizations from receiving future prenotifications if they release fixes or any other information about issues before they are public. Group membership may also change based on policy refinements for who may be included.
diff --git a/benchmarks/README.md b/benchmarks/README.md
index d6442a4fc3..98b3600d13 100644
--- a/benchmarks/README.md
+++ b/benchmarks/README.md
@@ -22,6 +22,25 @@ become available.
✅ |
wget https://huggingface.co/datasets/anon8231489123/ShareGPT_Vicuna_unfiltered/resolve/main/ShareGPT_V3_unfiltered_cleaned_split.json |
+
+ | ShareGPT4V (Image) |
+ ✅ |
+ ✅ |
+
+ wget https://huggingface.co/datasets/Lin-Chen/ShareGPT4V/blob/main/sharegpt4v_instruct_gpt4-vision_cap100k.json
+
+ Note that the images need to be downloaded separately. For example, to download COCO's 2017 Train images:
+ wget http://images.cocodataset.org/zips/train2017.zip
+ |
+
+
+ | ShareGPT4Video (Video) |
+ ✅ |
+ ✅ |
+
+ git clone https://huggingface.co/datasets/ShareGPT4Video/ShareGPT4Video
+ |
+
| BurstGPT |
✅ |
@@ -29,7 +48,7 @@ become available.
wget https://github.com/HPMLL/BurstGPT/releases/download/v1.1/BurstGPT_without_fails_2.csv |
- | Sonnet |
+ Sonnet (deprecated) |
✅ |
✅ |
Local file: benchmarks/sonnet.txt |
@@ -40,6 +59,18 @@ become available.
✅ |
synthetic |
+
+ | RandomMultiModal (Image/Video) |
+ 🟡 |
+ 🚧 |
+ synthetic |
+
+
+ | Prefix Repetition |
+ ✅ |
+ ✅ |
+ synthetic |
+
| HuggingFace-VisionArena |
✅ |
@@ -79,7 +110,12 @@ become available.
🚧: to be supported
-**Note**: HuggingFace dataset's `dataset-name` should be set to `hf`
+**Note**: HuggingFace dataset's `dataset-name` should be set to `hf`.
+For local `dataset-path`, please set `hf-name` to its Hugging Face ID like
+
+```bash
+--dataset-path /datasets/VisionArena-Chat/ --hf-name lmarena-ai/VisionArena-Chat
+```
## 🚀 Example - Online Benchmark
@@ -177,6 +213,7 @@ vllm serve Qwen/Qwen2-VL-7B-Instruct
```bash
vllm bench serve \
--backend openai-chat \
+ --endpoint-type openai-chat \
--model Qwen/Qwen2-VL-7B-Instruct \
--endpoint /v1/chat/completions \
--dataset-name hf \
@@ -213,6 +250,7 @@ vllm serve Qwen/Qwen2-VL-7B-Instruct
```bash
vllm bench serve \
--backend openai-chat \
+ --endpoint-type openai-chat \
--model Qwen/Qwen2-VL-7B-Instruct \
--endpoint /v1/chat/completions \
--dataset-name hf \
@@ -227,6 +265,7 @@ vllm bench serve \
```bash
vllm bench serve \
--backend openai-chat \
+ --endpoint-type openai-chat \
--model Qwen/Qwen2-VL-7B-Instruct \
--endpoint /v1/chat/completions \
--dataset-name hf \
@@ -581,6 +620,20 @@ python3 benchmarks/benchmark_prefix_caching.py \
--input-length-range 128:256
```
+### Prefix Repetition Dataset
+
+```bash
+vllm bench serve \
+ --backend openai \
+ --model meta-llama/Llama-2-7b-chat-hf \
+ --dataset-name prefix_repetition \
+ --num-prompts 100 \
+ --prefix-repetition-prefix-len 512 \
+ --prefix-repetition-suffix-len 128 \
+ --prefix-repetition-num-prefixes 5 \
+ --prefix-repetition-output-len 128
+```
+
## ⚡ Example - Request Prioritization Benchmark
@@ -616,3 +669,139 @@ python3 benchmarks/benchmark_prioritization.py \
```
+
+## 👁️ Example - Multi-Modal Benchmark
+
+
+Show more
+
+
+
+Benchmark the performance of multi-modal requests in vLLM.
+
+### Images (ShareGPT4V)
+
+Start vLLM:
+
+```bash
+python -m vllm.entrypoints.openai.api_server \
+ --model Qwen/Qwen2.5-VL-7B-Instruct \
+ --dtype bfloat16 \
+ --limit-mm-per-prompt '{"image": 1}' \
+ --allowed-local-media-path /path/to/sharegpt4v/images
+```
+
+Send requests with images:
+
+```bash
+python benchmarks/benchmark_serving.py \
+ --backend openai-chat \
+ --model Qwen/Qwen2.5-VL-7B-Instruct \
+ --dataset-name sharegpt \
+ --dataset-path /path/to/ShareGPT4V/sharegpt4v_instruct_gpt4-vision_cap100k.json \
+ --num-prompts 100 \
+ --save-result \
+ --result-dir ~/vllm_benchmark_results \
+ --save-detailed \
+ --endpoint /v1/chat/completion
+```
+
+### Videos (ShareGPT4Video)
+
+Start vLLM:
+
+```bash
+python -m vllm.entrypoints.openai.api_server \
+ --model Qwen/Qwen2.5-VL-7B-Instruct \
+ --dtype bfloat16 \
+ --limit-mm-per-prompt '{"video": 1}' \
+ --allowed-local-media-path /path/to/sharegpt4video/videos
+```
+
+Send requests with videos:
+
+```bash
+python benchmarks/benchmark_serving.py \
+ --backend openai-chat \
+ --model Qwen/Qwen2.5-VL-7B-Instruct \
+ --dataset-name sharegpt \
+ --dataset-path /path/to/ShareGPT4Video/llava_v1_5_mix665k_with_video_chatgpt72k_share4video28k.json \
+ --num-prompts 100 \
+ --save-result \
+ --result-dir ~/vllm_benchmark_results \
+ --save-detailed \
+ --endpoint /v1/chat/completion
+```
+
+### Synthetic Random Images (random-mm)
+
+Generate synthetic image inputs alongside random text prompts to stress-test vision models without external datasets.
+
+Notes:
+
+- Works only with online benchmark via the OpenAI backend (`--backend openai-chat`) and endpoint `/v1/chat/completions`.
+- Video sampling is not yet implemented.
+
+Start the server (example):
+
+```bash
+vllm serve Qwen/Qwen2.5-VL-3B-Instruct \
+ --dtype bfloat16 \
+ --max-model-len 16384 \
+ --limit-mm-per-prompt '{"image": 3, "video": 0}' \
+ --mm-processor-kwargs max_pixels=1003520
+```
+
+Benchmark. It is recommended to use the flag `--ignore-eos` to simulate real responses. You can set the size of the output via the arg `random-output-len`.
+
+Ex.1: Fixed number of items and a single image resolution, enforcing generation of approx 40 tokens:
+
+```bash
+vllm bench serve \
+ --backend openai-chat \
+ --model Qwen/Qwen2.5-VL-3B-Instruct \
+ --endpoint /v1/chat/completions \
+ --dataset-name random-mm \
+ --num-prompts 100 \
+ --max-concurrency 10 \
+ --random-prefix-len 25 \
+ --random-input-len 300 \
+ --random-output-len 40 \
+ --random-range-ratio 0.2 \
+ --random-mm-base-items-per-request 2 \
+ --random-mm-limit-mm-per-prompt '{"image": 3, "video": 0}' \
+ --random-mm-bucket-config '{(224, 224, 1): 1.0}' \
+ --request-rate inf \
+ --ignore-eos \
+ --seed 42
+```
+
+The number of items per request can be controlled by passing multiple image buckets:
+
+```bash
+ --random-mm-base-items-per-request 2 \
+ --random-mm-num-mm-items-range-ratio 0.5 \
+ --random-mm-limit-mm-per-prompt '{"image": 4, "video": 0}' \
+ --random-mm-bucket-config '{(256, 256, 1): 0.7, (720, 1280, 1): 0.3}' \
+```
+
+Flags specific to `random-mm`:
+
+- `--random-mm-base-items-per-request`: base number of multimodal items per request.
+- `--random-mm-num-mm-items-range-ratio`: vary item count uniformly in the closed integer range [floor(n·(1−r)), ceil(n·(1+r))]. Set r=0 to keep it fixed; r=1 allows 0 items.
+- `--random-mm-limit-mm-per-prompt`: per-modality hard caps, e.g. '{"image": 3, "video": 0}'.
+- `--random-mm-bucket-config`: dict mapping (H, W, T) → probability. Entries with probability 0 are removed; remaining probabilities are renormalized to sum to 1. Use T=1 for images. Set any T>1 for videos (video sampling not yet supported).
+
+Behavioral notes:
+
+- If the requested base item count cannot be satisfied under the provided per-prompt limits, the tool raises an error rather than silently clamping.
+
+How sampling works:
+
+- Determine per-request item count k by sampling uniformly from the integer range defined by `--random-mm-base-items-per-request` and `--random-mm-num-mm-items-range-ratio`, then clamp k to at most the sum of per-modality limits.
+- For each of the k items, sample a bucket (H, W, T) according to the normalized probabilities in `--random-mm-bucket-config`, while tracking how many items of each modality have been added.
+- If a modality (e.g., image) reaches its limit from `--random-mm-limit-mm-per-prompt`, all buckets of that modality are excluded and the remaining bucket probabilities are renormalized before continuing.
+This should be seen as an edge case, and if this behavior can be avoided by setting `--random-mm-limit-mm-per-prompt` to a large number. Note that this might result in errors due to engine config `--limit-mm-per-prompt`.
+- The resulting request contains synthetic image data in `multi_modal_data` (OpenAI Chat format). When `random-mm` is used with the OpenAI Chat backend, prompts remain text and MM content is attached via `multi_modal_data`.
+
+
diff --git a/benchmarks/auto_tune/README.md b/benchmarks/auto_tune/README.md
index 9aad51df6e..3aa988aac2 100644
--- a/benchmarks/auto_tune/README.md
+++ b/benchmarks/auto_tune/README.md
@@ -31,6 +31,12 @@ cd vllm
You must set the following variables at the top of the script before execution.
+ Note: You can also override the default values below via environment variables when running the script.
+
+```bash
+MODEL=meta-llama/Llama-3.3-70B-Instruct SYSTEM=TPU TP=8 DOWNLOAD_DIR='' INPUT_LEN=128 OUTPUT_LEN=2048 MAX_MODEL_LEN=2300 MIN_CACHE_HIT_PCT=0 MAX_LATENCY_ALLOWED_MS=100000000000 NUM_SEQS_LIST="128 256" NUM_BATCHED_TOKENS_LIST="1024 2048 4096" VLLM_LOGGING_LEVEL=DEBUG bash auto_tune.sh
+```
+
| Variable | Description | Example Value |
| --- | --- | --- |
| `BASE` | **Required.** The absolute path to the parent directory of your vLLM repository directory. | `"$HOME"` |
diff --git a/benchmarks/auto_tune/auto_tune.sh b/benchmarks/auto_tune/auto_tune.sh
index 82c20ffa65..ed3679b66f 100644
--- a/benchmarks/auto_tune/auto_tune.sh
+++ b/benchmarks/auto_tune/auto_tune.sh
@@ -5,25 +5,41 @@
TAG=$(date +"%Y_%m_%d_%H_%M")
SCRIPT_DIR=$( cd -- "$( dirname -- "${BASH_SOURCE[0]}" )" &> /dev/null && pwd )
-BASE="$SCRIPT_DIR/../../.."
-MODEL="meta-llama/Llama-3.1-8B-Instruct"
-SYSTEM="TPU"
-TP=1
-DOWNLOAD_DIR=""
-INPUT_LEN=4000
-OUTPUT_LEN=16
-MAX_MODEL_LEN=4096
-MIN_CACHE_HIT_PCT=0
-MAX_LATENCY_ALLOWED_MS=100000000000
-NUM_SEQS_LIST="128 256"
-NUM_BATCHED_TOKENS_LIST="512 1024 2048 4096"
+VLLM_LOGGING_LEVEL=${VLLM_LOGGING_LEVEL:-INFO}
+BASE=${BASE:-"$SCRIPT_DIR/../../.."}
+MODEL=${MODEL:-"meta-llama/Llama-3.1-8B-Instruct"}
+SYSTEM=${SYSTEM:-"TPU"}
+TP=${TP:-1}
+DOWNLOAD_DIR=${DOWNLOAD_DIR:-""}
+INPUT_LEN=${INPUT_LEN:-4000}
+OUTPUT_LEN=${OUTPUT_LEN:-16}
+MAX_MODEL_LEN=${MAX_MODEL_LEN:-4096}
+MIN_CACHE_HIT_PCT=${MIN_CACHE_HIT_PCT:-0}
+MAX_LATENCY_ALLOWED_MS=${MAX_LATENCY_ALLOWED_MS:-100000000000}
+NUM_SEQS_LIST=${NUM_SEQS_LIST:-"128 256"}
+NUM_BATCHED_TOKENS_LIST=${NUM_BATCHED_TOKENS_LIST:-"512 1024 2048 4096"}
LOG_FOLDER="$BASE/auto-benchmark/$TAG"
RESULT="$LOG_FOLDER/result.txt"
PROFILE_PATH="$LOG_FOLDER/profile"
-echo "result file: $RESULT"
-echo "model: $MODEL"
+echo "====================== AUTO TUNE PARAMETERS ===================="
+echo "SCRIPT_DIR=$SCRIPT_DIR"
+echo "BASE=$BASE"
+echo "MODEL=$MODEL"
+echo "SYSTEM=$SYSTEM"
+echo "TP=$TP"
+echo "DOWNLOAD_DIR=$DOWNLOAD_DIR"
+echo "INPUT_LEN=$INPUT_LEN"
+echo "OUTPUT_LEN=$OUTPUT_LEN"
+echo "MAX_MODEL_LEN=$MAX_MODEL_LEN"
+echo "MIN_CACHE_HIT_PCT=$MIN_CACHE_HIT_PCT"
+echo "MAX_LATENCY_ALLOWED_MS=$MAX_LATENCY_ALLOWED_MS"
+echo "NUM_SEQS_LIST=$NUM_SEQS_LIST"
+echo "NUM_BATCHED_TOKENS_LIST=$NUM_BATCHED_TOKENS_LIST"
+echo "VLLM_LOGGING_LEVEL=$VLLM_LOGGING_LEVEL"
+echo "RESULT_FILE=$RESULT"
+echo "====================== AUTO TUNEPARAMETERS ===================="
rm -rf $LOG_FOLDER
rm -rf $PROFILE_PATH
@@ -213,7 +229,7 @@ run_benchmark() {
pkill -if vllm
sleep 10
- printf '=%.0s' $(seq 1 20)
+ echo "===================="
return 0
}
diff --git a/benchmarks/backend_request_func.py b/benchmarks/backend_request_func.py
index c7229dbb8e..ba7c733be0 100644
--- a/benchmarks/backend_request_func.py
+++ b/benchmarks/backend_request_func.py
@@ -31,9 +31,10 @@ class RequestFuncInput:
model_name: Optional[str] = None
logprobs: Optional[int] = None
extra_body: Optional[dict] = None
- multi_modal_content: Optional[dict] = None
+ multi_modal_content: Optional[dict | list[dict]] = None
ignore_eos: bool = False
language: Optional[str] = None
+ request_id: Optional[str] = None
@dataclass
@@ -71,6 +72,9 @@ async def async_request_tgi(
"inputs": request_func_input.prompt,
"parameters": params,
}
+ headers = None
+ if request_func_input.request_id:
+ headers = {"x-request-id": request_func_input.request_id}
output = RequestFuncOutput()
output.prompt_len = request_func_input.prompt_len
if request_func_input.ignore_eos:
@@ -82,7 +86,9 @@ async def async_request_tgi(
st = time.perf_counter()
most_recent_timestamp = st
try:
- async with session.post(url=api_url, json=payload) as response:
+ async with session.post(
+ url=api_url, json=payload, headers=headers
+ ) as response:
if response.status == 200:
async for chunk_bytes in response.content:
chunk_bytes = chunk_bytes.strip()
@@ -145,6 +151,9 @@ async def async_request_trt_llm(
}
if request_func_input.ignore_eos:
payload["min_length"] = request_func_input.output_len
+ headers = None
+ if request_func_input.request_id:
+ headers = {"x-request-id": request_func_input.request_id}
output = RequestFuncOutput()
output.prompt_len = request_func_input.prompt_len
@@ -152,7 +161,9 @@ async def async_request_trt_llm(
st = time.perf_counter()
most_recent_timestamp = st
try:
- async with session.post(url=api_url, json=payload) as response:
+ async with session.post(
+ url=api_url, json=payload, headers=headers
+ ) as response:
if response.status == 200:
async for chunk_bytes in response.content:
chunk_bytes = chunk_bytes.strip()
@@ -211,6 +222,8 @@ async def async_request_deepspeed_mii(
"top_p": 1.0,
}
headers = {"Authorization": f"Bearer {os.environ.get('OPENAI_API_KEY')}"}
+ if request_func_input.request_id:
+ headers["x-request-id"] = request_func_input.request_id
output = RequestFuncOutput()
output.prompt_len = request_func_input.prompt_len
@@ -283,6 +296,8 @@ async def async_request_openai_completions(
if request_func_input.extra_body:
payload.update(request_func_input.extra_body)
headers = {"Authorization": f"Bearer {os.environ.get('OPENAI_API_KEY')}"}
+ if request_func_input.request_id:
+ headers["x-request-id"] = request_func_input.request_id
output = RequestFuncOutput()
output.prompt_len = request_func_input.prompt_len
@@ -364,7 +379,15 @@ async def async_request_openai_chat_completions(
) as session:
content = [{"type": "text", "text": request_func_input.prompt}]
if request_func_input.multi_modal_content:
- content.append(request_func_input.multi_modal_content)
+ mm_content = request_func_input.multi_modal_content
+ if isinstance(mm_content, list):
+ content.extend(mm_content)
+ elif isinstance(mm_content, dict):
+ content.append(mm_content)
+ else:
+ raise TypeError(
+ "multi_modal_content must be a dict or list[dict] for openai-chat"
+ )
payload = {
"model": request_func_input.model_name
if request_func_input.model_name
@@ -387,6 +410,8 @@ async def async_request_openai_chat_completions(
"Content-Type": "application/json",
"Authorization": f"Bearer {os.environ.get('OPENAI_API_KEY')}",
}
+ if request_func_input.request_id:
+ headers["x-request-id"] = request_func_input.request_id
output = RequestFuncOutput()
output.prompt_len = request_func_input.prompt_len
@@ -483,6 +508,8 @@ async def async_request_openai_audio(
headers = {
"Authorization": f"Bearer {os.environ.get('OPENAI_API_KEY')}",
}
+ if request_func_input.request_id:
+ headers["x-request-id"] = request_func_input.request_id
# Send audio file
def to_bytes(y, sr):
@@ -491,7 +518,10 @@ async def async_request_openai_audio(
buffer.seek(0)
return buffer
- with to_bytes(*request_func_input.multi_modal_content["audio"]) as f:
+ mm_audio = request_func_input.multi_modal_content
+ if not isinstance(mm_audio, dict) or "audio" not in mm_audio:
+ raise TypeError("multi_modal_content must be a dict containing 'audio'")
+ with to_bytes(*mm_audio["audio"]) as f:
form = aiohttp.FormData()
form.add_field("file", f, content_type="audio/wav")
for key, value in payload.items():
diff --git a/benchmarks/benchmark_block_pool.py b/benchmarks/benchmark_block_pool.py
new file mode 100644
index 0000000000..eae8d9927e
--- /dev/null
+++ b/benchmarks/benchmark_block_pool.py
@@ -0,0 +1,74 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import gc
+
+from tabulate import tabulate
+
+from benchmark_utils import TimeCollector
+from vllm.utils import FlexibleArgumentParser
+from vllm.v1.core.block_pool import BlockPool
+
+
+def main(args):
+ rows = []
+ for allocate_block in args.allocate_blocks:
+ # Enforce a GC collect ahead to minimize the impact among runs
+ gc.collect()
+ block_pool = BlockPool(num_gpu_blocks=args.num_gpu_blocks, enable_caching=True)
+
+ get_blocks_times = TimeCollector(TimeCollector.US)
+ free_blocks_times = TimeCollector(TimeCollector.US)
+ for _ in range(args.num_iteration):
+ with get_blocks_times:
+ blocks = block_pool.get_new_blocks(allocate_block)
+ with free_blocks_times:
+ block_pool.free_blocks(blocks)
+
+ rows.append(
+ [get_blocks_times.cnt, args.num_gpu_blocks, allocate_block]
+ + get_blocks_times.dump_avg_max()
+ + free_blocks_times.dump_avg_max()
+ )
+
+ print(
+ tabulate(
+ rows,
+ headers=[
+ "Iterations",
+ "Total\nBlocks",
+ "Allocated\nBlocks",
+ "Get Blocks\nAvg (us)",
+ "Get Blocks\nMax (us)",
+ "Free Blocks\nAvg (us)",
+ "Free Blocks\nMax (us)",
+ ],
+ tablefmt="grid",
+ floatfmt=".3f",
+ )
+ )
+
+
+def invoke_main() -> None:
+ parser = FlexibleArgumentParser(
+ description="Benchmark the performance of BlockPool for KV Cache."
+ )
+ parser.add_argument("--num-gpu-blocks", type=int, default=100000)
+ parser.add_argument(
+ "--num-iteration",
+ type=int,
+ default=1000,
+ help="Number of iterations to run to stabilize final data readings",
+ )
+ parser.add_argument(
+ "--allocate-blocks",
+ type=int,
+ nargs="*",
+ default=[10, 50, 100, 500, 1000],
+ help="Number of blocks to allocate",
+ )
+ args = parser.parse_args()
+ main(args)
+
+
+if __name__ == "__main__":
+ invoke_main() # pragma: no cover
diff --git a/benchmarks/benchmark_dataset.py b/benchmarks/benchmark_dataset.py
index 1ad6cef7a9..64ffa62c04 100644
--- a/benchmarks/benchmark_dataset.py
+++ b/benchmarks/benchmark_dataset.py
@@ -19,6 +19,7 @@ import logging
import random
from abc import ABC, abstractmethod
from collections.abc import Mapping
+from copy import deepcopy
from dataclasses import dataclass
from functools import cache
from io import BytesIO
@@ -52,8 +53,9 @@ class SampleRequest:
prompt: Union[str, Any]
prompt_len: int
expected_output_len: int
- multi_modal_data: Optional[Union[MultiModalDataDict, dict]] = None
+ multi_modal_data: Optional[Union[MultiModalDataDict, dict, list[dict]]] = None
lora_request: Optional[LoRARequest] = None
+ request_id: Optional[str] = None
# -----------------------------------------------------------------------------
@@ -155,7 +157,10 @@ class BenchmarkDataset(ABC):
@abstractmethod
def sample(
- self, tokenizer: PreTrainedTokenizerBase, num_requests: int
+ self,
+ tokenizer: PreTrainedTokenizerBase,
+ num_requests: int,
+ request_id_prefix: str = "",
) -> list[SampleRequest]:
"""
Abstract method to generate sample requests from the dataset.
@@ -167,6 +172,7 @@ class BenchmarkDataset(ABC):
tokenizer (PreTrainedTokenizerBase): The tokenizer to be used
for processing the dataset's text.
num_requests (int): The number of sample requests to generate.
+ request_id_prefix (str) The prefix of request_id.
Returns:
list[SampleRequest]: A list of sample requests generated from the
@@ -175,7 +181,10 @@ class BenchmarkDataset(ABC):
raise NotImplementedError("sample must be implemented in subclasses.")
def maybe_oversample_requests(
- self, requests: list[SampleRequest], num_requests: int
+ self,
+ requests: list[SampleRequest],
+ num_requests: int,
+ request_id_prefix: str = "",
) -> None:
"""
Oversamples the list of requests if its size is less than the desired
@@ -183,11 +192,18 @@ class BenchmarkDataset(ABC):
Args:
requests (List[SampleRequest]): The current list of sampled
- requests. num_requests (int): The target number of requests.
+ requests.
+ num_requests (int): The target number of requests.
+ request_id_prefix (str) The prefix of the request ids.
"""
if len(requests) < num_requests:
random.seed(self.random_seed)
- additional = random.choices(requests, k=num_requests - len(requests))
+ additional = deepcopy(
+ random.choices(requests, k=num_requests - len(requests))
+ )
+ for i in range(len(additional)):
+ req = additional[i]
+ req.request_id = request_id_prefix + str(len(requests) + i)
requests.extend(additional)
logger.info("Oversampled requests to reach %d total samples.", num_requests)
@@ -277,6 +293,41 @@ def process_image(image: Any) -> Mapping[str, Any]:
)
+def process_video(video: Any) -> Mapping[str, Any]:
+ """
+ Process a single video input and return a multimedia content dictionary.
+
+ Supports the following input types:
+
+ 1. Dictionary with raw video bytes: - Expects a dict with a 'bytes' key
+ containing raw video data.
+
+ 2. String input: - Treats the string as a URL or local file path. -
+ Prepends "file://" if the string doesn't start with "http://" or
+ "file://". - Returns a dictionary with the image URL.
+
+ Raises:
+ ValueError: If the input is not a supported type.
+ """
+ if isinstance(video, dict) and "bytes" in video:
+ video_bytes = video["bytes"]
+ video_base64 = base64.b64encode(video_bytes).decode("utf-8")
+ return {
+ "type": "video_url",
+ "video_url": {"url": f"data:video/mp4;base64,{video_base64}"},
+ }
+
+ if isinstance(video, str):
+ video_url = (
+ video if video.startswith(("http://", "file://")) else f"file://{video}"
+ )
+ return {"type": "video_url", "video_url": {"url": video_url}}
+
+ raise ValueError(
+ f"Invalid video input {video}. Must be a string of local path/remote url, or a dictionary with raw video bytes in the form of `{{'bytes': raw_video_bytes}}`." # noqa: E501
+ )
+
+
# -----------------------------------------------------------------------------
# Random Dataset Implementation (Synthetic Data)
# -----------------------------------------------------------------------------
@@ -303,6 +354,7 @@ class RandomDataset(BenchmarkDataset):
range_ratio: float = DEFAULT_RANGE_RATIO,
input_len: int = DEFAULT_INPUT_LEN,
output_len: int = DEFAULT_OUTPUT_LEN,
+ request_id_prefix: str = "",
**kwargs,
) -> list[SampleRequest]:
# Enforce range_ratio < 1
@@ -351,7 +403,7 @@ class RandomDataset(BenchmarkDataset):
# [6880, 6881] -> ['Ġcalls', 'here'] ->
# [1650, 939, 486] -> ['Ġcall', 'sh', 'ere']
# To avoid uncontrolled change of the prompt length,
- # the encoded sequence is truncated before being decode again.
+ # the encoded sequence is truncated before being decoded again.
total_input_len = prefix_len + int(input_lens[i])
re_encoded_sequence = tokenizer.encode(prompt, add_special_tokens=False)[
:total_input_len
@@ -363,8 +415,10 @@ class RandomDataset(BenchmarkDataset):
prompt=prompt,
prompt_len=total_input_len,
expected_output_len=int(output_lens[i]),
+ request_id=request_id_prefix + str(i),
)
)
+
return requests
@@ -406,9 +460,11 @@ class ShareGPTDataset(BenchmarkDataset):
max_loras: Optional[int] = None,
output_len: Optional[int] = None,
enable_multimodal_chat: bool = False,
+ request_id_prefix: str = "",
**kwargs,
) -> list:
samples: list = []
+ ind = 0
for entry in self.data:
if len(samples) >= num_requests:
break
@@ -430,17 +486,26 @@ class ShareGPTDataset(BenchmarkDataset):
skip_min_output_len_check=output_len is not None,
):
continue
+ if image_path := entry.get("image"):
+ mm_content = process_image(image_path)
+ elif video_path := entry.get("video"):
+ mm_content = process_video(video_path)
+ else:
+ mm_content = None
if enable_multimodal_chat:
- prompt = self.apply_multimodal_chat_transformation(prompt, None)
+ prompt = self.apply_multimodal_chat_transformation(prompt, mm_content)
samples.append(
SampleRequest(
prompt=prompt,
prompt_len=prompt_len,
expected_output_len=new_output_len,
lora_request=lora_request,
+ multi_modal_data=mm_content,
+ request_id=request_id_prefix + str(ind),
)
)
- self.maybe_oversample_requests(samples, num_requests)
+ ind += 1
+ self.maybe_oversample_requests(samples, num_requests, request_id_prefix)
return samples
@@ -506,10 +571,11 @@ class CustomDataset(BenchmarkDataset):
output_len: Optional[int] = None,
enable_multimodal_chat: bool = False,
skip_chat_template: bool = False,
+ request_id_prefix: str = "",
**kwargs,
) -> list:
sampled_requests = []
- for item in self.data:
+ for i, item in enumerate(self.data):
if len(sampled_requests) >= num_requests:
break
prompt = item["prompt"]
@@ -528,9 +594,12 @@ class CustomDataset(BenchmarkDataset):
prompt=prompt,
prompt_len=prompt_len,
expected_output_len=output_len,
+ request_id=request_id_prefix + str(i),
)
)
- self.maybe_oversample_requests(sampled_requests, num_requests)
+ self.maybe_oversample_requests(
+ sampled_requests, num_requests, request_id_prefix
+ )
return sampled_requests
@@ -572,6 +641,7 @@ class SonnetDataset(BenchmarkDataset):
input_len: int = DEFAULT_INPUT_LEN,
output_len: int = DEFAULT_OUTPUT_LEN,
return_prompt_formatted: bool = False,
+ request_id_prefix: str = "",
**kwargs,
) -> list:
# Calculate average token length for a poem line.
@@ -597,6 +667,7 @@ class SonnetDataset(BenchmarkDataset):
prefix_lines = self.data[:num_prefix_lines]
samples = []
+ ind = 0
while len(samples) < num_requests:
extra_lines = random.choices(
self.data, k=num_input_lines - num_prefix_lines
@@ -607,14 +678,17 @@ class SonnetDataset(BenchmarkDataset):
msg, add_generation_prompt=True, tokenize=False
)
prompt_len = len(tokenizer(prompt_formatted).input_ids)
+
if prompt_len <= input_len:
samples.append(
SampleRequest(
prompt=prompt_formatted if return_prompt_formatted else prompt,
prompt_len=prompt_len,
expected_output_len=output_len,
+ request_id=request_id_prefix + str(ind),
)
)
+ ind += 1
return samples
@@ -666,6 +740,7 @@ class BurstGPTDataset(BenchmarkDataset):
num_requests: int,
max_loras: Optional[int] = None,
lora_path: Optional[str] = None,
+ request_id_prefix: str = "",
**kwargs,
) -> list[SampleRequest]:
samples = []
@@ -687,6 +762,7 @@ class BurstGPTDataset(BenchmarkDataset):
prompt_len=input_len,
expected_output_len=output_len,
lora_request=lora_req,
+ request_id=request_id_prefix + str(i),
)
)
return samples
@@ -746,12 +822,14 @@ class ConversationDataset(HuggingFaceDataset):
num_requests: int,
output_len: Optional[int] = None,
enable_multimodal_chat: bool = False,
+ request_id_prefix: str = "",
**kwargs,
) -> list:
# Filter examples with at least 2 conversations
filtered_data = self.data.filter(lambda x: len(x["conversations"]) >= 2)
sampled_requests = []
dynamic_output = output_len is None
+ ind = 0
for item in filtered_data:
if len(sampled_requests) >= num_requests:
@@ -779,9 +857,13 @@ class ConversationDataset(HuggingFaceDataset):
prompt_len=prompt_len,
expected_output_len=output_len,
multi_modal_data=mm_content,
+ request_id=request_id_prefix + str(ind),
)
)
- self.maybe_oversample_requests(sampled_requests, num_requests)
+ ind += 1
+ self.maybe_oversample_requests(
+ sampled_requests, num_requests, request_id_prefix
+ )
return sampled_requests
@@ -808,11 +890,12 @@ class VisionArenaDataset(HuggingFaceDataset):
num_requests: int,
output_len: Optional[int] = None,
enable_multimodal_chat: bool = False,
+ request_id_prefix: str = "",
**kwargs,
) -> list:
output_len = output_len if output_len is not None else self.DEFAULT_OUTPUT_LEN
sampled_requests = []
- for item in self.data:
+ for i, item in enumerate(self.data):
if len(sampled_requests) >= num_requests:
break
parser_fn = self.SUPPORTED_DATASET_PATHS.get(self.dataset_path)
@@ -832,9 +915,12 @@ class VisionArenaDataset(HuggingFaceDataset):
prompt_len=prompt_len,
expected_output_len=output_len,
multi_modal_data=mm_content,
+ request_id=request_id_prefix + str(i),
)
)
- self.maybe_oversample_requests(sampled_requests, num_requests)
+ self.maybe_oversample_requests(
+ sampled_requests, num_requests, request_id_prefix
+ )
return sampled_requests
@@ -864,15 +950,18 @@ class InstructCoderDataset(HuggingFaceDataset):
num_requests: int,
output_len: Optional[int] = None,
enable_multimodal_chat: bool = False,
+ request_id_prefix: str = "",
**kwargs,
) -> list:
output_len = output_len if output_len is not None else self.DEFAULT_OUTPUT_LEN
sampled_requests = []
- for item in self.data:
+ for i, item in enumerate(self.data):
if len(sampled_requests) >= num_requests:
break
- prompt = f"{item['input']}\n\n{item['instruction']} Just output \
- the code, do not include any explanation."
+ prompt = (
+ f"{item['input']}\n\n{item['instruction']} Just output "
+ "the code, do not include any explanation."
+ )
# apply template
prompt = tokenizer.apply_chat_template(
@@ -886,9 +975,12 @@ class InstructCoderDataset(HuggingFaceDataset):
prompt=prompt,
prompt_len=prompt_len,
expected_output_len=output_len,
+ request_id=request_id_prefix + str(i),
)
)
- self.maybe_oversample_requests(sampled_requests, num_requests)
+ self.maybe_oversample_requests(
+ sampled_requests, num_requests, request_id_prefix
+ )
return sampled_requests
@@ -918,12 +1010,13 @@ class MTBenchDataset(HuggingFaceDataset):
num_requests: int,
output_len: Optional[int] = None,
enable_multimodal_chat: bool = False,
+ request_id_prefix: str = "",
**kwargs,
) -> list:
output_len = output_len if output_len is not None else self.DEFAULT_OUTPUT_LEN
sampled_requests = []
- for item in self.data:
+ for i, item in enumerate(self.data):
if len(sampled_requests) >= num_requests:
break
prompt = item["turns"][0]
@@ -941,9 +1034,12 @@ class MTBenchDataset(HuggingFaceDataset):
prompt=prompt,
prompt_len=prompt_len,
expected_output_len=output_len,
+ request_id=request_id_prefix + str(i),
)
)
- self.maybe_oversample_requests(sampled_requests, num_requests)
+ self.maybe_oversample_requests(
+ sampled_requests, num_requests, request_id_prefix
+ )
return sampled_requests
@@ -968,10 +1064,12 @@ class AIMODataset(HuggingFaceDataset):
tokenizer: PreTrainedTokenizerBase,
num_requests: int,
output_len: Optional[int] = None,
+ request_id_prefix: str = "",
**kwargs,
) -> list:
sampled_requests = []
dynamic_output = output_len is None
+ ind = 0
for item in self.data:
if len(sampled_requests) >= num_requests:
@@ -994,9 +1092,13 @@ class AIMODataset(HuggingFaceDataset):
prompt_len=prompt_len,
expected_output_len=output_len,
multi_modal_data=None,
+ request_id=request_id_prefix + str(ind),
)
)
- self.maybe_oversample_requests(sampled_requests, num_requests)
+ ind += 1
+ self.maybe_oversample_requests(
+ sampled_requests, num_requests, request_id_prefix
+ )
return sampled_requests
@@ -1066,12 +1168,18 @@ class NextEditPredictionDataset(HuggingFaceDataset):
"zed-industries/zeta": _format_zeta_prompt,
}
- def sample(self, tokenizer: PreTrainedTokenizerBase, num_requests: int, **kwargs):
+ def sample(
+ self,
+ tokenizer: PreTrainedTokenizerBase,
+ num_requests: int,
+ request_id_prefix: str = "",
+ **kwargs,
+ ):
formatting_prompt_func = self.MAPPING_PROMPT_FUNCS.get(self.dataset_path)
if formatting_prompt_func is None:
raise ValueError(f"Unsupported dataset path: {self.dataset_path}")
samples = []
- for sample in self.data:
+ for i, sample in enumerate(self.data):
sample = formatting_prompt_func(sample)
samples.append(
SampleRequest(
@@ -1080,11 +1188,12 @@ class NextEditPredictionDataset(HuggingFaceDataset):
expected_output_len=len(
tokenizer(sample["expected_output"]).input_ids
),
+ request_id=request_id_prefix + str(i),
)
)
if len(samples) >= num_requests:
break
- self.maybe_oversample_requests(samples, num_requests)
+ self.maybe_oversample_requests(samples, num_requests, request_id_prefix)
return samples
@@ -1133,6 +1242,7 @@ class ASRDataset(HuggingFaceDataset):
tokenizer: PreTrainedTokenizerBase,
num_requests: int,
output_len: Optional[int] = None,
+ request_id_prefix: str = "",
**kwargs,
) -> list:
import librosa
@@ -1142,6 +1252,7 @@ class ASRDataset(HuggingFaceDataset):
prompt_len = len(tokenizer(prompt).input_ids)
sampled_requests = []
skipped = 0
+ ind = 0
for item in self.data:
if len(sampled_requests) >= num_requests:
break
@@ -1160,8 +1271,10 @@ class ASRDataset(HuggingFaceDataset):
prompt_len=prompt_len,
expected_output_len=output_len,
multi_modal_data=mm_content,
+ request_id=request_id_prefix + str(ind),
)
)
+ ind += 1
if skipped:
logger.warning(
"%d samples discarded from dataset due to"
@@ -1169,5 +1282,7 @@ class ASRDataset(HuggingFaceDataset):
" what Whisper supports.",
skipped,
)
- self.maybe_oversample_requests(sampled_requests, num_requests)
+ self.maybe_oversample_requests(
+ sampled_requests, num_requests, request_id_prefix
+ )
return sampled_requests
diff --git a/benchmarks/benchmark_ngram_proposer.py b/benchmarks/benchmark_ngram_proposer.py
new file mode 100644
index 0000000000..11833fa1b3
--- /dev/null
+++ b/benchmarks/benchmark_ngram_proposer.py
@@ -0,0 +1,112 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import gc
+
+import numpy as np
+from tabulate import tabulate
+
+from benchmark_utils import TimeCollector
+from vllm.config import ModelConfig, SpeculativeConfig, VllmConfig
+from vllm.utils import FlexibleArgumentParser
+from vllm.v1.spec_decode.ngram_proposer import NgramProposer
+
+
+def main(args):
+ rows = []
+ for max_ngram in args.max_ngram:
+ collector = TimeCollector(TimeCollector.US)
+
+ model_config = ModelConfig(
+ model="facebook/opt-125m",
+ task="generate",
+ max_model_len=args.num_token + args.num_spec_token,
+ tokenizer="facebook/opt-125m",
+ tokenizer_mode="auto",
+ dtype="auto",
+ seed=None,
+ trust_remote_code=False,
+ )
+ proposer = NgramProposer(
+ vllm_config=VllmConfig(
+ model_config=model_config,
+ speculative_config=SpeculativeConfig(
+ prompt_lookup_min=args.min_ngram,
+ prompt_lookup_max=max_ngram,
+ num_speculative_tokens=args.num_spec_token,
+ method="ngram",
+ ),
+ )
+ )
+
+ # Warm up
+ proposer.propose(np.random.randint(0, 20, (args.num_token,)))
+
+ gc.collect()
+ for _ in range(args.num_iteration):
+ tokens = np.random.randint(0, 20, (args.num_req, args.num_token))
+ with collector:
+ for i in range(args.num_req):
+ proposer.propose(tokens[i, :])
+ rows.append(
+ [args.num_req, args.num_token, args.min_ngram, max_ngram]
+ + collector.dump_avg_max()
+ )
+
+ print(
+ tabulate(
+ rows,
+ headers=[
+ "# Request",
+ "# Token",
+ "Min Ngram",
+ "Max Ngram",
+ "Avg (us)",
+ "Max (us)",
+ ],
+ tablefmt="grid",
+ floatfmt=".3f",
+ )
+ )
+
+
+def invoke_main() -> None:
+ parser = FlexibleArgumentParser(
+ description="Benchmark the performance of N-gram speculative decode drafting"
+ )
+ parser.add_argument(
+ "--num-iteration",
+ type=int,
+ default=100,
+ help="Number of iterations to run to stabilize final data readings",
+ )
+ parser.add_argument(
+ "--num-req", type=int, default=128, help="Number of requests in the batch"
+ )
+ parser.add_argument(
+ "--num-token", type=int, default=1500, help="Number of tokens for each request"
+ )
+ parser.add_argument(
+ "--min-ngram",
+ type=int,
+ default=3,
+ help="Minimum n-gram to match",
+ )
+ parser.add_argument(
+ "--max-ngram",
+ type=int,
+ nargs="*",
+ default=[5, 7, 10, 15, 20],
+ help="Maximum n-gram to match",
+ )
+ parser.add_argument(
+ "--num-spec-token",
+ type=int,
+ default=3,
+ help="Number of speculative tokens to generate",
+ )
+ args = parser.parse_args()
+ main(args)
+
+
+if __name__ == "__main__":
+ invoke_main() # pragma: no cover
diff --git a/benchmarks/benchmark_serving.py b/benchmarks/benchmark_serving.py
index 93b72211eb..934df05efa 100644
--- a/benchmarks/benchmark_serving.py
+++ b/benchmarks/benchmark_serving.py
@@ -263,7 +263,14 @@ async def benchmark(
input_requests[0].multi_modal_data,
)
- assert test_mm_content is None or isinstance(test_mm_content, dict)
+ assert (
+ test_mm_content is None
+ or isinstance(test_mm_content, dict)
+ or (
+ isinstance(test_mm_content, list)
+ and all(isinstance(item, dict) for item in test_mm_content)
+ )
+ ), "multi_modal_data must be a dict or list[dict]"
test_input = RequestFuncInput(
model=model_id,
model_name=model_name,
@@ -368,11 +375,12 @@ async def benchmark(
rps_change_events.append({"rps": rps_val, "timestamp": timestamp})
last_int_rps = current_int_rps
- prompt, prompt_len, output_len, mm_content = (
+ prompt, prompt_len, output_len, mm_content, request_id = (
request.prompt,
request.prompt_len,
request.expected_output_len,
request.multi_modal_data,
+ request.request_id,
)
req_model_id, req_model_name = model_id, model_name
if lora_modules:
@@ -390,6 +398,7 @@ async def benchmark(
multi_modal_content=mm_content,
ignore_eos=ignore_eos,
extra_body=extra_body,
+ request_id=request_id,
)
task = limited_request_func(request_func_input=request_func_input, pbar=pbar)
tasks.append(asyncio.create_task(task))
@@ -658,6 +667,7 @@ def main(args: argparse.Namespace):
tokenizer=tokenizer,
output_len=args.custom_output_len,
skip_chat_template=args.custom_skip_chat_template,
+ request_id_prefix=args.request_id_prefix,
)
elif args.dataset_name == "sonnet":
@@ -671,6 +681,7 @@ def main(args: argparse.Namespace):
prefix_len=args.sonnet_prefix_len,
tokenizer=tokenizer,
return_prompt_formatted=False,
+ request_id_prefix=args.request_id_prefix,
)
else:
assert tokenizer.chat_template or tokenizer.default_chat_template, (
@@ -683,6 +694,7 @@ def main(args: argparse.Namespace):
prefix_len=args.sonnet_prefix_len,
tokenizer=tokenizer,
return_prompt_formatted=True,
+ request_id_prefix=args.request_id_prefix,
)
elif args.dataset_name == "hf":
@@ -744,6 +756,7 @@ def main(args: argparse.Namespace):
num_requests=args.num_prompts,
tokenizer=tokenizer,
output_len=args.hf_output_len,
+ request_id_prefix=args.request_id_prefix,
)
else:
@@ -755,10 +768,15 @@ def main(args: argparse.Namespace):
tokenizer=tokenizer,
num_requests=args.num_prompts,
output_len=args.sharegpt_output_len,
+ request_id_prefix=args.request_id_prefix,
),
"burstgpt": lambda: BurstGPTDataset(
random_seed=args.seed, dataset_path=args.dataset_path
- ).sample(tokenizer=tokenizer, num_requests=args.num_prompts),
+ ).sample(
+ tokenizer=tokenizer,
+ num_requests=args.num_prompts,
+ request_id_prefix=args.request_id_prefix,
+ ),
"random": lambda: RandomDataset(dataset_path=args.dataset_path).sample(
tokenizer=tokenizer,
num_requests=args.num_prompts,
@@ -766,6 +784,7 @@ def main(args: argparse.Namespace):
input_len=args.random_input_len,
output_len=args.random_output_len,
range_ratio=args.random_range_ratio,
+ request_id_prefix=args.request_id_prefix,
),
}
@@ -1085,7 +1104,7 @@ def create_argument_parser():
"--percentile-metrics",
type=str,
default="ttft,tpot,itl",
- help="Comma-separated list of selected metrics to report percentils. "
+ help="Comma-separated list of selected metrics to report percentiles. "
"This argument specifies the metrics to report percentiles. "
'Allowed metric names are "ttft", "tpot", "itl", "e2el". '
'Default value is "ttft,tpot,itl".',
@@ -1111,6 +1130,13 @@ def create_argument_parser():
"goodput, refer to DistServe paper: https://arxiv.org/pdf/2401.09670 "
"and the blog: https://hao-ai-lab.github.io/blogs/distserve",
)
+ parser.add_argument(
+ "--request-id-prefix",
+ type=str,
+ required=False,
+ default="benchmark-serving",
+ help="Specify the prefix of request id.",
+ )
# group for dataset specific arguments
custom_group = parser.add_argument_group("custom dataset options")
diff --git a/benchmarks/benchmark_serving_structured_output.py b/benchmarks/benchmark_serving_structured_output.py
index ca6843a72a..4aae755eb4 100644
--- a/benchmarks/benchmark_serving_structured_output.py
+++ b/benchmarks/benchmark_serving_structured_output.py
@@ -998,7 +998,7 @@ def create_argument_parser():
"--percentile-metrics",
type=str,
default="ttft,tpot,itl",
- help="Comma-separated list of selected metrics to report percentils. "
+ help="Comma-separated list of selected metrics to report percentiles. "
"This argument specifies the metrics to report percentiles. "
'Allowed metric names are "ttft", "tpot", "itl", "e2el". '
'Default value is "ttft,tpot,itl".',
diff --git a/benchmarks/benchmark_throughput.py b/benchmarks/benchmark_throughput.py
index c51b579686..34a525f00d 100644
--- a/benchmarks/benchmark_throughput.py
+++ b/benchmarks/benchmark_throughput.py
@@ -96,7 +96,6 @@ def run_vllm(
end = time.perf_counter()
else:
assert lora_requests is None, "BeamSearch API does not support LoRA"
- prompts = [request.prompt for request in requests]
# output_len should be the same for all requests.
output_len = requests[0].expected_output_len
for request in requests:
@@ -597,8 +596,8 @@ def validate_args(args):
# https://github.com/vllm-project/vllm/issues/16222
if args.data_parallel_size > 1:
raise ValueError(
- "Data parallel is not supported in offline benchmark, \
- please use benchmark serving instead"
+ "Data parallel is not supported in offline benchmark, "
+ "please use benchmark serving instead"
)
@@ -720,7 +719,7 @@ def create_argument_parser():
"[length * (1 - range_ratio), length * (1 + range_ratio)].",
)
- # hf dtaset
+ # hf dataset
parser.add_argument(
"--hf-subset", type=str, default=None, help="Subset of the HF dataset."
)
diff --git a/benchmarks/benchmark_utils.py b/benchmarks/benchmark_utils.py
index 283f938df5..98624abdf4 100644
--- a/benchmarks/benchmark_utils.py
+++ b/benchmarks/benchmark_utils.py
@@ -1,11 +1,12 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
import argparse
import json
import math
import os
-from typing import Any
+import time
+from types import TracebackType
+from typing import Any, Optional, Union
def convert_to_pytorch_benchmark_format(
@@ -72,3 +73,53 @@ def write_to_json(filename: str, records: list) -> None:
cls=InfEncoder,
default=lambda o: f"<{type(o).__name__} object is not JSON serializable>",
)
+
+
+# Collect time and generate time metrics
+#
+# Example Usage:
+# collector = TimeCollector(TimeCollector.US)
+# for _ in range(total_iteration):
+# with collector:
+# ...
+# collector.dump_avg_max()
+class TimeCollector:
+ NS: int = 1
+ US: int = NS * 1000
+ MS: int = US * 1000
+ S: int = MS * 1000
+
+ def __init__(self, scale: int) -> None:
+ self.cnt: int = 0
+ self._sum: int = 0
+ self._max: Optional[int] = None
+ self.scale = scale
+ self.start_time: int = time.monotonic_ns()
+
+ def collect(self, v: int) -> None:
+ self.cnt += 1
+ self._sum += v
+ if self._max is None:
+ self._max = v
+ else:
+ self._max = max(self._max, v)
+
+ def avg(self) -> Union[float, str]:
+ return self._sum * 1.0 / self.cnt / self.scale if self.cnt > 0 else "N/A"
+
+ def max(self) -> Union[float, str]:
+ return self._max / self.scale if self._max else "N/A"
+
+ def dump_avg_max(self) -> list[Union[float, str]]:
+ return [self.avg(), self.max()]
+
+ def __enter__(self) -> None:
+ self.start_time = time.monotonic_ns()
+
+ def __exit__(
+ self,
+ exc_type: Optional[type[BaseException]],
+ exc_value: Optional[BaseException],
+ exc_traceback: Optional[TracebackType],
+ ) -> None:
+ self.collect(time.monotonic_ns() - self.start_time)
diff --git a/benchmarks/disagg_benchmarks/disagg_overhead_benchmark.sh b/benchmarks/disagg_benchmarks/disagg_overhead_benchmark.sh
index 92f97ffabe..2c72941cf7 100644
--- a/benchmarks/disagg_benchmarks/disagg_overhead_benchmark.sh
+++ b/benchmarks/disagg_benchmarks/disagg_overhead_benchmark.sh
@@ -62,7 +62,7 @@ benchmark() {
--max-model-len 10000 \
--gpu-memory-utilization 0.6 \
--kv-transfer-config \
- '{"kv_connector":"PyNcclConnector","kv_role":"kv_producer","kv_rank":0,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
+ '{"kv_connector":"P2pNcclConnector","kv_role":"kv_producer","kv_rank":0,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
CUDA_VISIBLE_DEVICES=1 python3 \
@@ -72,7 +72,7 @@ benchmark() {
--max-model-len 10000 \
--gpu-memory-utilization 0.6 \
--kv-transfer-config \
- '{"kv_connector":"PyNcclConnector","kv_role":"kv_consumer","kv_rank":1,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
+ '{"kv_connector":"P2pNcclConnector","kv_role":"kv_consumer","kv_rank":1,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
wait_for_server 8100
wait_for_server 8200
diff --git a/benchmarks/disagg_benchmarks/disagg_performance_benchmark.sh b/benchmarks/disagg_benchmarks/disagg_performance_benchmark.sh
index af2bcba3ea..0bbf7cd2b1 100644
--- a/benchmarks/disagg_benchmarks/disagg_performance_benchmark.sh
+++ b/benchmarks/disagg_benchmarks/disagg_performance_benchmark.sh
@@ -69,7 +69,7 @@ launch_disagg_prefill() {
--max-model-len 10000 \
--gpu-memory-utilization 0.6 \
--kv-transfer-config \
- '{"kv_connector":"PyNcclConnector","kv_role":"kv_producer","kv_rank":0,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
+ '{"kv_connector":"P2pNcclConnector","kv_role":"kv_producer","kv_rank":0,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
CUDA_VISIBLE_DEVICES=1 python3 \
-m vllm.entrypoints.openai.api_server \
@@ -78,7 +78,7 @@ launch_disagg_prefill() {
--max-model-len 10000 \
--gpu-memory-utilization 0.6 \
--kv-transfer-config \
- '{"kv_connector":"PyNcclConnector","kv_role":"kv_consumer","kv_rank":1,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
+ '{"kv_connector":"P2pNcclConnector","kv_role":"kv_consumer","kv_rank":1,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
wait_for_server 8100
wait_for_server 8200
diff --git a/benchmarks/disagg_benchmarks/disagg_prefill_proxy_server.py b/benchmarks/disagg_benchmarks/disagg_prefill_proxy_server.py
index f62d8102e2..904f805349 100644
--- a/benchmarks/disagg_benchmarks/disagg_prefill_proxy_server.py
+++ b/benchmarks/disagg_benchmarks/disagg_prefill_proxy_server.py
@@ -1,63 +1,199 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import argparse
+import asyncio
+import logging
import os
import aiohttp
-from quart import Quart, make_response, request
+from quart import Quart, Response, make_response, request
+from rate_limiter import RateLimiter
+from request_queue import RequestQueue
-AIOHTTP_TIMEOUT = aiohttp.ClientTimeout(total=6 * 60 * 60)
-
-app = Quart(__name__)
+# Configure logging
+logging.basicConfig(level=logging.INFO)
+logger = logging.getLogger(__name__)
-async def forward_request(url, data):
- async with aiohttp.ClientSession(timeout=AIOHTTP_TIMEOUT) as session:
+def parse_args():
+ """parse command line arguments"""
+ parser = argparse.ArgumentParser(description="vLLM P/D disaggregation proxy server")
+
+ # Add args
+ parser.add_argument(
+ "--timeout",
+ type=float,
+ default=300,
+ help="Timeout for backend service requests in seconds (default: 300)",
+ )
+ parser.add_argument(
+ "--max-concurrent",
+ type=int,
+ default=100,
+ help="Maximum concurrent requests to backend services (default: 100)",
+ )
+ parser.add_argument(
+ "--queue-size",
+ type=int,
+ default=500,
+ help="Maximum number of requests in the queue (default: 500)",
+ )
+ parser.add_argument(
+ "--rate-limit",
+ type=int,
+ default=40,
+ help="Maximum requests per second (default: 40)",
+ )
+ parser.add_argument(
+ "--port",
+ type=int,
+ default=8000,
+ help="Port to run the server on (default: 8000)",
+ )
+ parser.add_argument(
+ "--prefill-url",
+ type=str,
+ default="http://localhost:8100/v1/completions",
+ help="Prefill service endpoint URL",
+ )
+ parser.add_argument(
+ "--decode-url",
+ type=str,
+ default="http://localhost:8200/v1/completions",
+ help="Decode service endpoint URL",
+ )
+
+ return parser.parse_args()
+
+
+def main():
+ """parse command line arguments"""
+ args = parse_args()
+
+ # Initialize configuration using command line parameters
+ AIOHTTP_TIMEOUT = aiohttp.ClientTimeout(total=args.timeout)
+ MAX_CONCURRENT_REQUESTS = args.max_concurrent
+ REQUEST_QUEUE_SIZE = args.queue_size
+ RATE_LIMIT = args.rate_limit
+ PREFILL_SERVICE_URL = args.prefill_url
+ DECODE_SERVICE_URL = args.decode_url
+ PORT = args.port
+
+ app = Quart(__name__)
+
+ # Initialize the rate limiter and request queue
+ rate_limiter = RateLimiter(RATE_LIMIT)
+ request_queue = RequestQueue(MAX_CONCURRENT_REQUESTS, REQUEST_QUEUE_SIZE)
+
+ # Attach the configuration object to the application instance
+ app.config.update(
+ {
+ "AIOHTTP_TIMEOUT": AIOHTTP_TIMEOUT,
+ "rate_limiter": rate_limiter,
+ "request_queue": request_queue,
+ "PREFILL_SERVICE_URL": PREFILL_SERVICE_URL,
+ "DECODE_SERVICE_URL": DECODE_SERVICE_URL,
+ }
+ )
+
+ # Start queue processing on app startup
+ @app.before_serving
+ async def startup():
+ """Start request processing task when app starts serving"""
+ asyncio.create_task(request_queue.process())
+
+ async def forward_request(url, data):
+ """Forward request to backend service with rate limiting and error handling"""
headers = {"Authorization": f"Bearer {os.environ.get('OPENAI_API_KEY')}"}
- async with session.post(url=url, json=data, headers=headers) as response:
- if response.status == 200:
- # if response.headers.get('Transfer-Encoding') == 'chunked':
- if True:
- async for chunk_bytes in response.content.iter_chunked(1024):
- yield chunk_bytes
- else:
- content = await response.read()
- yield content
-
-@app.route("/v1/completions", methods=["POST"])
-async def handle_request():
- try:
- original_request_data = await request.get_json()
-
- prefill_request = original_request_data.copy()
- # change max_tokens = 1 to let it only do prefill
- prefill_request["max_tokens"] = 1
-
- # finish prefill
- async for _ in forward_request(
- "http://localhost:8100/v1/completions", prefill_request
+ # Use rate limiter as context manager
+ async with (
+ rate_limiter,
+ aiohttp.ClientSession(timeout=AIOHTTP_TIMEOUT) as session,
):
- continue
+ try:
+ async with session.post(
+ url=url, json=data, headers=headers
+ ) as response:
+ if response.status == 200:
+ # Stream response chunks
+ async for chunk_bytes in response.content.iter_chunked(1024):
+ yield chunk_bytes
+ else:
+ # Handle backend service errors
+ error_text = await response.text()
+ logger.error(
+ "Backend service error: %s - %s",
+ response.status,
+ error_text,
+ )
+ yield b'{"error": "Backend service error"}'
+ except aiohttp.ClientError as e:
+ # Handle connection errors
+ logger.error("Connection error to %s: %s", url, str(e))
+ yield b'{"error": "Service unavailable"}'
+ except asyncio.TimeoutError:
+ # Handle timeout errors
+ logger.error("Timeout connecting to %s", url)
+ yield b'{"error": "Service timeout"}'
- # return decode
- generator = forward_request(
- "http://localhost:8200/v1/completions", original_request_data
- )
- response = await make_response(generator)
- response.timeout = None
+ async def process_request():
+ """Process a single request through prefill and decode stages"""
+ try:
+ original_request_data = await request.get_json()
- return response
+ # Create prefill request (max_tokens=1)
+ prefill_request = original_request_data.copy()
+ prefill_request["max_tokens"] = 1
- except Exception as e:
- import sys
- import traceback
+ # Execute prefill stage
+ async for _ in forward_request(PREFILL_SERVICE_URL, prefill_request):
+ continue
- exc_info = sys.exc_info()
- print("Error occurred in disagg prefill proxy server")
- print(e)
- print("".join(traceback.format_exception(*exc_info)))
+ # Execute decode stage and stream response
+ generator = forward_request(DECODE_SERVICE_URL, original_request_data)
+ response = await make_response(generator)
+ response.timeout = None # Disable timeout for streaming response
+ return response
+
+ except Exception:
+ logger.exception("Error processing request")
+ return Response(
+ response=b'{"error": "Internal server error"}',
+ status=500,
+ content_type="application/json",
+ )
+
+ @app.route("/v1/completions", methods=["POST"])
+ async def handle_request():
+ """Handle incoming API requests with concurrency and rate limiting"""
+ # Create task for request processing
+ task = asyncio.create_task(process_request())
+
+ # Enqueue request or reject if queue is full
+ if not await request_queue.enqueue(task):
+ return Response(
+ response=b'{"error": "Server busy, try again later"}',
+ status=503,
+ content_type="application/json",
+ )
+
+ try:
+ # Return the response from the processing task
+ return await task
+ except asyncio.CancelledError:
+ # Handle task cancellation (timeout or queue full)
+ logger.warning("Request cancelled due to timeout or queue full")
+ return Response(
+ response=b'{"error": "Request cancelled"}',
+ status=503,
+ content_type="application/json",
+ )
+
+ # Start the Quart server with host can be set to 0.0.0.0
+ app.run(port=PORT)
if __name__ == "__main__":
- app.run(port=8000)
+ main()
diff --git a/benchmarks/disagg_benchmarks/rate_limiter.py b/benchmarks/disagg_benchmarks/rate_limiter.py
new file mode 100644
index 0000000000..87ac8cb6ab
--- /dev/null
+++ b/benchmarks/disagg_benchmarks/rate_limiter.py
@@ -0,0 +1,45 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import asyncio
+import time
+
+
+class RateLimiter:
+ """Token bucket rate limiter implementation"""
+
+ def __init__(self, rate_limit):
+ self.rate_limit = rate_limit # Requests per second
+ self.num_available_tokens = rate_limit # Available tokens
+ self.last_refill = time.monotonic() # Last token refill time
+ self.lock = asyncio.Lock() # Synchronization lock
+
+ async def acquire(self):
+ """Acquire a token from the rate limiter"""
+ while True:
+ async with self.lock:
+ current_time = time.monotonic()
+ elapsed = current_time - self.last_refill
+
+ # Refill num_available_tokens if more than 1 second has passed
+ if elapsed > 1.0:
+ self.num_available_tokens = self.rate_limit
+ self.last_refill = current_time
+
+ # Check if num_available_tokens are available
+ if self.num_available_tokens > 0:
+ self.num_available_tokens -= 1
+ return True
+
+ # Calculate wait time if no num_available_tokens available
+ wait_time = 1.0 - elapsed
+ await asyncio.sleep(wait_time)
+
+ async def __aenter__(self):
+ """Enter async context manager - acquire token"""
+ await self.acquire()
+ return self
+
+ async def __aexit__(self, exc_type, exc_value, traceback):
+ """Exit async context manager - no cleanup needed"""
+ pass
diff --git a/benchmarks/disagg_benchmarks/request_queue.py b/benchmarks/disagg_benchmarks/request_queue.py
new file mode 100644
index 0000000000..410bcb9560
--- /dev/null
+++ b/benchmarks/disagg_benchmarks/request_queue.py
@@ -0,0 +1,39 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import asyncio
+from collections import deque
+
+
+class RequestQueue:
+ """Request queue manager with concurrency control"""
+
+ def __init__(self, max_concurrent, max_queue_size):
+ # Maximum concurrent requests
+ self.max_concurrent = max_concurrent
+ self.max_queue_size = max_queue_size # Maximum queue size
+ # Concurrency control
+ self.semaphore = asyncio.Semaphore(max_concurrent)
+ self.queue = deque() # Request queue
+ self.queue_size = 0 # Current queue size
+ self.lock = asyncio.Lock() # Sync queue Lock
+
+ async def enqueue(self, task):
+ """Add a request task to the queue"""
+ async with self.lock:
+ if self.queue_size >= self.max_queue_size:
+ return False
+
+ self.queue.append(task)
+ self.queue_size += 1
+ return True
+
+ async def process(self):
+ """Process queued requests using semaphore for concurrency control"""
+ while True:
+ if self.queue:
+ async with self.semaphore, self.lock:
+ task = self.queue.popleft()
+ self.queue_size -= 1
+ await task
+ await asyncio.sleep(0.01) # Yield control to event loop
diff --git a/benchmarks/kernels/bench_block_fp8_gemm.py b/benchmarks/kernels/bench_block_fp8_gemm.py
new file mode 100644
index 0000000000..9663503e9b
--- /dev/null
+++ b/benchmarks/kernels/bench_block_fp8_gemm.py
@@ -0,0 +1,114 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import torch
+
+from vllm.model_executor.layers.quantization.utils.fp8_utils import (
+ w8a8_block_fp8_matmul,
+)
+from vllm.platforms import current_platform
+from vllm.triton_utils import triton as vllm_triton
+
+assert current_platform.is_cuda(), (
+ "Only support benchmarking w8a8 block fp8 kernel on CUDA device."
+)
+
+# DeepSeek-V3 weight shapes
+DEEPSEEK_V3_SHAPES = [
+ (512 + 64, 7168),
+ (2112, 7168),
+ ((128 + 64) * 128, 7168),
+ (128 * (128 + 128), 512),
+ (7168, 16384),
+ (7168, 18432),
+ (18432 * 2, 7168),
+ (24576, 1536),
+ (12288, 7168),
+ (4096, 7168),
+ (7168, 2048),
+]
+
+
+def build_w8a8_block_fp8_runner(M, N, K, block_size, device):
+ """Build runner function for w8a8 block fp8 matmul."""
+ factor_for_scale = 1e-2
+
+ fp8_info = torch.finfo(torch.float8_e4m3fn)
+ fp8_max, fp8_min = fp8_info.max, fp8_info.min
+
+ # Create random FP8 tensors
+ A_fp32 = (torch.rand(M, K, dtype=torch.float32, device=device) - 0.5) * 2 * fp8_max
+ A = A_fp32.clamp(min=fp8_min, max=fp8_max).to(torch.float8_e4m3fn)
+
+ B_fp32 = (torch.rand(N, K, dtype=torch.float32, device=device) - 0.5) * 2 * fp8_max
+ B = B_fp32.clamp(min=fp8_min, max=fp8_max).to(torch.float8_e4m3fn)
+
+ # Create scales
+ block_n, block_k = block_size[0], block_size[1]
+ n_tiles = (N + block_n - 1) // block_n
+ k_tiles = (K + block_k - 1) // block_k
+
+ As = torch.rand(M, k_tiles, dtype=torch.float32, device=device) * factor_for_scale
+ Bs = (
+ torch.rand(n_tiles, k_tiles, dtype=torch.float32, device=device)
+ * factor_for_scale
+ )
+
+ def run():
+ return w8a8_block_fp8_matmul(A, B, As, Bs, block_size, torch.bfloat16)
+
+ return run
+
+
+@vllm_triton.testing.perf_report(
+ vllm_triton.testing.Benchmark(
+ x_names=["batch_size"],
+ x_vals=[1, 16, 64, 128, 256, 512, 1024, 2048, 4096, 8192, 16384],
+ x_log=False,
+ line_arg="provider",
+ line_vals=["torch-bf16", "w8a8-block-fp8"],
+ line_names=["torch-bf16", "w8a8-block-fp8"],
+ ylabel="TFLOP/s (larger is better)",
+ plot_name="BF16 vs W8A8 Block FP8 GEMMs",
+ args={},
+ )
+)
+def benchmark_tflops(batch_size, provider, N, K, block_size=(128, 128)):
+ M = batch_size
+ device = "cuda"
+
+ quantiles = [0.5, 0.2, 0.8]
+
+ if provider == "torch-bf16":
+ a = torch.randn((M, K), device=device, dtype=torch.bfloat16)
+ b = torch.randn((N, K), device=device, dtype=torch.bfloat16)
+ ms, min_ms, max_ms = vllm_triton.testing.do_bench_cudagraph(
+ lambda: torch.nn.functional.linear(a, b), quantiles=quantiles
+ )
+ else: # w8a8-block-fp8
+ run_w8a8 = build_w8a8_block_fp8_runner(M, N, K, block_size, device)
+ ms, min_ms, max_ms = vllm_triton.testing.do_bench_cudagraph(
+ lambda: run_w8a8(), quantiles=quantiles
+ )
+
+ to_tflops = lambda t_ms: (2 * M * N * K) * 1e-12 / (t_ms * 1e-3)
+ return to_tflops(ms), to_tflops(max_ms), to_tflops(min_ms)
+
+
+if __name__ == "__main__":
+ block_size = (128, 128)
+
+ for N, K in DEEPSEEK_V3_SHAPES:
+ print(f"\nBenchmarking DeepSeek-V3, N={N} K={K}")
+
+ print(f"TFLOP/s comparison (block_size={block_size}):")
+ benchmark_tflops.run(
+ print_data=True,
+ # show_plots=False,
+ # save_path=f"bench_w8a8_block_fp8_tflops_n{N}_k{K}",
+ N=N,
+ K=K,
+ block_size=block_size,
+ )
+
+ print("\nBenchmark finished!")
diff --git a/benchmarks/kernels/benchmark_activation.py b/benchmarks/kernels/benchmark_activation.py
new file mode 100644
index 0000000000..93edbcc939
--- /dev/null
+++ b/benchmarks/kernels/benchmark_activation.py
@@ -0,0 +1,104 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+# benchmark custom activation op performance
+import itertools
+
+import torch
+
+import vllm.model_executor.layers.activation # noqa F401
+from vllm.model_executor.custom_op import CustomOp
+from vllm.platforms import current_platform
+from vllm.triton_utils import triton
+from vllm.utils import STR_DTYPE_TO_TORCH_DTYPE, FlexibleArgumentParser
+
+batch_size_range = [1, 16, 32, 64, 128]
+seq_len_range = [1, 16, 64, 128, 256, 512, 1024, 2048, 4096]
+intermediate_size = [3072, 9728, 12288]
+configs = list(itertools.product(batch_size_range, seq_len_range, intermediate_size))
+
+
+def benchmark_activation(
+ batch_size: int,
+ seq_len: int,
+ intermediate_size: int,
+ provider: str,
+ func_name: str,
+ dtype: torch.dtype,
+):
+ device = "cuda"
+ num_tokens = batch_size * seq_len
+ dim = intermediate_size
+ current_platform.seed_everything(42)
+ torch.set_default_device(device)
+
+ if func_name == "gelu_and_mul":
+ layer = CustomOp.op_registry[func_name](approximate="none")
+ elif func_name == "gelu_and_mul_tanh":
+ layer = CustomOp.op_registry["gelu_and_mul"](approximate="tanh")
+ elif func_name == "fatrelu_and_mul":
+ threshold = 0.5
+ layer = CustomOp.op_registry[func_name](threshold)
+ else:
+ layer = CustomOp.op_registry[func_name]()
+
+ x = torch.randn(num_tokens, dim, dtype=dtype, device=device)
+ compiled_layer = torch.compile(layer.forward_native)
+
+ if provider == "custom":
+ fn = lambda: layer(x)
+ elif provider == "compiled":
+ fn = lambda: compiled_layer(x)
+
+ ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
+ fn, quantiles=[0.5, 0.2, 0.8]
+ )
+ return ms, max_ms, min_ms
+
+
+if __name__ == "__main__":
+ parser = FlexibleArgumentParser(description="Benchmark the custom activation op.")
+ parser.add_argument(
+ "--func-name",
+ type=str,
+ choices=[
+ "mul_and_silu",
+ "silu_and_mul",
+ "gelu_and_mul",
+ "gelu_and_mul_tanh",
+ "fatrelu_and_mul",
+ "swigluoai_and_mul",
+ "gelu_new",
+ "gelu_fast",
+ "quick_gelu",
+ ],
+ default="silu_and_mul",
+ )
+ parser.add_argument(
+ "--dtype", type=str, choices=["half", "bfloat16", "float"], default="bfloat16"
+ )
+ args = parser.parse_args()
+ assert args
+
+ func_name = args.func_name
+ dtype = STR_DTYPE_TO_TORCH_DTYPE[args.dtype]
+
+ perf_report = triton.testing.perf_report(
+ triton.testing.Benchmark(
+ x_names=["batch_size", "seq_len", "intermediate_size"],
+ x_vals=configs,
+ line_arg="provider",
+ line_vals=["custom", "compiled"],
+ line_names=["Custom OP", "Compiled"],
+ styles=[("blue", "-"), ("green", "-")],
+ ylabel="ms",
+ plot_name=f"{func_name}-op-performance",
+ args={},
+ )
+ )
+
+ perf_report(
+ lambda batch_size, seq_len, intermediate_size, provider: benchmark_activation(
+ batch_size, seq_len, intermediate_size, provider, func_name, dtype
+ )
+ ).run(print_data=True)
diff --git a/benchmarks/kernels/benchmark_aqlm.py b/benchmarks/kernels/benchmark_aqlm.py
deleted file mode 100644
index 42de062b08..0000000000
--- a/benchmarks/kernels/benchmark_aqlm.py
+++ /dev/null
@@ -1,345 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-import os
-import sys
-from typing import Optional
-
-import torch
-import torch.nn.functional as F
-
-from vllm import _custom_ops as ops
-from vllm.model_executor.layers.quantization.aqlm import (
- dequantize_weight,
- generic_dequantize_gemm,
- get_int_dtype,
- optimized_dequantize_gemm,
-)
-from vllm.utils import FlexibleArgumentParser
-
-os.environ["CUDA_VISIBLE_DEVICES"] = "0"
-
-
-def torch_mult(
- # [..., in_features]
- input: torch.Tensor,
- weights: torch.Tensor,
- # [num_out_groups, 1, 1, 1]
- scales: torch.Tensor,
-) -> torch.Tensor:
- output = F.linear(input, weights)
- return output
-
-
-def dequant_out_scale(
- # [..., in_features]
- input: torch.Tensor,
- # [num_out_groups, num_in_groups, num_codebooks]
- codes: torch.IntTensor,
- # [num_codebooks, codebook_size, out_group_size, in_group_size]
- codebooks: torch.Tensor,
- # [num_out_groups, 1, 1, 1]
- scales: torch.Tensor,
- output_partition_sizes: torch.IntTensor,
- bias: Optional[torch.Tensor],
-) -> torch.Tensor:
- weights = ops.aqlm_dequant(codes, codebooks, output_partition_sizes)
-
- if bias is None:
- output = F.linear(input, weights, bias)
- orig_shape = output.shape
- flattened_output = output.view(-1, output.size(-1))
- f_scales = scales.view(-1, scales.shape[0])
- b_scales = f_scales.expand(flattened_output.shape[0], -1)
- flattened_output *= b_scales
- return flattened_output.view(orig_shape)
- else:
- b_scales = scales.view(scales.shape[:-3] + (-1,)).expand(-1, weights.shape[1])
- weights *= b_scales
- return F.linear(input, weights, bias)
-
-
-def dequant_weight_scale(
- # [..., in_features]
- input: torch.Tensor,
- # [num_out_groups, num_in_groups, num_codebooks]
- codes: torch.IntTensor,
- # [num_codebooks, codebook_size, out_group_size, in_group_size]
- codebooks: torch.Tensor,
- # [num_out_groups, 1, 1, 1]
- scales: torch.Tensor,
- output_partition_sizes: torch.IntTensor,
- bias: Optional[torch.Tensor],
-) -> torch.Tensor:
- weights = ops.aqlm_dequant(codes, codebooks, output_partition_sizes)
-
- b_scales = scales.view(scales.shape[:-3] + (-1,)).expand(-1, weights.shape[1])
- weights *= b_scales
- return F.linear(input, weights, bias)
-
-
-def dequant_no_scale(
- # [..., in_features]
- input: torch.Tensor,
- # [num_out_groups, num_in_groups, num_codebooks]
- codes: torch.IntTensor,
- # [num_codebooks, codebook_size, out_group_size, in_group_size]
- codebooks: torch.Tensor,
- # [num_out_groups, 1, 1, 1]
- scales: torch.Tensor,
- output_partition_sizes: torch.IntTensor,
- bias: Optional[torch.Tensor],
-) -> torch.Tensor:
- weights = ops.aqlm_dequant(codes, codebooks, output_partition_sizes)
-
- return F.linear(input, weights, bias)
-
-
-# Compare the optimized 1x16 and 2x8 cuda decompression/dequant kernels against
-# the generic pytorch version.
-# Just visual comparison.
-def dequant_test(k: int, parts: torch.Tensor, nbooks: int, bits: int) -> None:
- n = int(parts.sum().item())
-
- device = torch.device("cuda:0")
-
- code_range = (1 << bits) // 2
- ingroups = 8
-
- codes = torch.randint(
- -code_range,
- code_range,
- size=(n, k // ingroups, nbooks),
- dtype=get_int_dtype(bits),
- device=device,
- )
-
- codebooks = torch.randn(
- size=(parts.shape[0] * nbooks, 1 << bits, 1, 8),
- dtype=torch.float16,
- device=device,
- )
-
- count = 0
- for index in range(16):
- for i in range(8):
- for book in range(nbooks):
- codebooks[book, index, 0, i] = count * (10**book)
- count += 1
-
- print("codes shape", codes.shape)
-
- for i in range(16):
- for book in range(nbooks):
- codes[0, i, book] = i
- codes[0, -i, book] = i
-
- weights = dequantize_weight(codes, codebooks, None)
- weights2 = ops.aqlm_dequant(codes, codebooks, parts)
-
- print("weights shape:", weights.shape)
- print("weights2 shape:", weights2.shape)
-
- print("weights are:", weights)
- print("weights2 are:", weights2)
-
- print("first 128 weights are", weights[0, 0:128].to(torch.int32))
- print("first 128 weights2 are:", weights2[0, 0:128].to(torch.int32))
-
- print("last 128 weights are", weights[0, -128:])
- print("last 128 weights2 are:", weights2[0, -128:])
-
-
-def main():
- parser = FlexibleArgumentParser(description="Benchmark aqlm performance.")
-
- # Add arguments
- parser.add_argument(
- "--nbooks", type=int, default=1, help="Number of codebooks (default: 1)"
- )
- parser.add_argument(
- "--bits",
- type=int,
- default=16,
- help="Number of bits per code element (default: 16)",
- )
- parser.add_argument(
- "--test",
- type=bool,
- default=False,
- help="Run the decompression/dequant tester rather than benchmarking "
- "(default: False)",
- )
-
- # Parse the arguments
- args = parser.parse_args()
-
- # Extract values
- nbooks = args.nbooks
- bits = args.bits
-
- if args.test:
- dequant_test(4096, torch.tensor((4096,)), nbooks, bits)
- return
-
- # Otherwise, benchmark.
- methods = [
- ops.aqlm_gemm,
- dequant_out_scale,
- generic_dequantize_gemm,
- optimized_dequantize_gemm,
- dequant_weight_scale,
- torch_mult,
- dequant_no_scale,
- ]
-
- filename = f"./aqlm_benchmark_{nbooks}x{bits}.csv"
- print(f"writing benchmarks to file {filename}")
- with open(filename, "w") as f:
- sys.stdout = f
-
- print("m | k | n | n parts", end="")
- for method in methods:
- print(f" | {method.__name__.replace('_', ' ')} (µs)", end="")
- print("")
-
- # These are reasonable prefill sizes.
- ksandpartions = (
- (4096, (4096, 4096, 4096)),
- (4096, (4096,)),
- (4096, (11008, 11008)),
- (11008, (4096,)),
- )
-
- # reasonable ranges for m.
- for m in [
- 1,
- 2,
- 4,
- 8,
- 10,
- 12,
- 14,
- 16,
- 24,
- 32,
- 48,
- 52,
- 56,
- 64,
- 96,
- 112,
- 128,
- 256,
- 512,
- 1024,
- 1536,
- 2048,
- 3072,
- 4096,
- ]:
- print(f"{m}", file=sys.__stdout__)
- for ksp in ksandpartions:
- run_grid(m, ksp[0], torch.tensor(ksp[1]), nbooks, bits, methods)
-
- sys.stdout = sys.__stdout__
-
-
-def run_grid(m: int, k: int, parts: torch.Tensor, nbooks: int, bits: int, methods):
- # I didn't see visible improvements from increasing these, but feel free :)
- num_warmup_trials = 1
- num_trials = 1
-
- num_calls = 100
-
- # warmup.
- for method in methods:
- for _ in range(num_warmup_trials):
- run_timing(
- num_calls=num_calls,
- m=m,
- k=k,
- parts=parts,
- nbooks=nbooks,
- bits=bits,
- method=method,
- )
-
- n = parts.sum().item()
- print(f"{m} | {k} | {n} | {parts.tolist()}", end="")
-
- for method in methods:
- best_time_us = 1e20
- for _ in range(num_trials):
- kernel_dur_ms = run_timing(
- num_calls=num_calls,
- m=m,
- k=k,
- parts=parts,
- nbooks=nbooks,
- bits=bits,
- method=method,
- )
-
- kernel_dur_us = 1000 * kernel_dur_ms
-
- if kernel_dur_us < best_time_us:
- best_time_us = kernel_dur_us
-
- print(f" | {kernel_dur_us:.0f}", end="")
-
- print("")
-
-
-def run_timing(
- num_calls: int, m: int, k: int, parts: torch.Tensor, nbooks: int, bits: int, method
-) -> float:
- n = int(parts.sum().item())
-
- device = torch.device("cuda:0")
-
- input = torch.randn((1, m, k), dtype=torch.float16, device=device)
-
- code_range = (1 << bits) // 2
- ingroups = 8
-
- codes = torch.randint(
- -code_range,
- code_range,
- size=(n, k // ingroups, nbooks),
- dtype=get_int_dtype(bits),
- device=device,
- )
-
- codebooks = torch.randn(
- size=(parts.shape[0] * nbooks, 1 << bits, 1, 8),
- dtype=torch.float16,
- device=device,
- )
-
- scales = torch.randn(size=(n, 1, 1, 1), dtype=torch.float16, device=device)
-
- # for comparison to just a pytorch mult.
- weights = torch.randn((n, k), dtype=torch.float16, device=device)
-
- start_event = torch.cuda.Event(enable_timing=True)
- end_event = torch.cuda.Event(enable_timing=True)
-
- start_event.record()
-
- if method is torch_mult:
- for i in range(num_calls):
- torch_mult(input, weights, scales)
- else:
- for i in range(num_calls):
- method(input, codes, codebooks, scales, parts, None)
-
- end_event.record()
- end_event.synchronize()
-
- dur_ms = start_event.elapsed_time(end_event) / num_calls
- return dur_ms
-
-
-if __name__ == "__main__":
- sys.exit(main())
diff --git a/benchmarks/kernels/benchmark_bitblas.py b/benchmarks/kernels/benchmark_bitblas.py
index 97ee060341..66b44c27d6 100644
--- a/benchmarks/kernels/benchmark_bitblas.py
+++ b/benchmarks/kernels/benchmark_bitblas.py
@@ -3,6 +3,8 @@
# Copyright (c) Microsoft Corporation.
# Licensed under the MIT License.
+from packaging import version
+
from vllm.model_executor.layers.quantization.utils.bitblas_utils import (
MINIMUM_BITBLAS_VERSION,
)
@@ -10,7 +12,7 @@ from vllm.model_executor.layers.quantization.utils.bitblas_utils import (
try:
import bitblas
- if bitblas.__version__ < MINIMUM_BITBLAS_VERSION:
+ if version.parse(bitblas.__version__) < version.parse(MINIMUM_BITBLAS_VERSION):
raise ImportError(
"bitblas version is wrong. Please "
f"install bitblas>={MINIMUM_BITBLAS_VERSION}"
diff --git a/benchmarks/kernels/benchmark_grouped_gemm_cutlass.py b/benchmarks/kernels/benchmark_grouped_gemm_cutlass.py
index 1d4e730f99..a6b42406b5 100644
--- a/benchmarks/kernels/benchmark_grouped_gemm_cutlass.py
+++ b/benchmarks/kernels/benchmark_grouped_gemm_cutlass.py
@@ -80,6 +80,11 @@ def bench_run(
a, score, topk, renormalize=False
)
+ ab_strides1 = torch.full((num_experts,), k, device="cuda", dtype=torch.int64)
+ ab_strides2 = torch.full((num_experts,), n, device="cuda", dtype=torch.int64)
+ c_strides1 = torch.full((num_experts,), 2 * n, device="cuda", dtype=torch.int64)
+ c_strides2 = torch.full((num_experts,), k, device="cuda", dtype=torch.int64)
+
def run_triton_moe(
a: torch.Tensor,
w1: torch.Tensor,
@@ -111,6 +116,10 @@ def bench_run(
w2: torch.Tensor,
w1_scale: torch.Tensor,
w2_scale: torch.Tensor,
+ ab_strides1: torch.Tensor,
+ ab_strides2: torch.Tensor,
+ c_strides1: torch.Tensor,
+ c_strides2: torch.Tensor,
topk_weights: torch.Tensor,
topk_ids: torch.Tensor,
per_act_token: bool,
@@ -125,6 +134,10 @@ def bench_run(
topk_ids,
w1_scale,
w2_scale,
+ ab_strides1,
+ ab_strides2,
+ c_strides1,
+ c_strides2,
per_act_token,
a1_scale=None,
)
@@ -136,6 +149,10 @@ def bench_run(
w2_q: torch.Tensor,
w1_scale: torch.Tensor,
w2_scale: torch.Tensor,
+ ab_strides1: torch.Tensor,
+ ab_strides2: torch.Tensor,
+ c_strides1: torch.Tensor,
+ c_strides2: torch.Tensor,
topk_weights: torch.Tensor,
topk_ids: torch.Tensor,
):
@@ -150,6 +167,10 @@ def bench_run(
topk_ids,
w1_scale,
w2_scale,
+ ab_strides1,
+ ab_strides2,
+ c_strides1,
+ c_strides2,
per_act_token,
a1_scale=None,
)
@@ -194,6 +215,10 @@ def bench_run(
w2_q,
w1_scale,
w2_scale,
+ ab_strides1,
+ ab_strides2,
+ c_strides1,
+ c_strides2,
topk_weights,
topk_ids,
)
@@ -231,6 +256,10 @@ def bench_run(
"w1_scale": w1_scale,
"w2_scale": w2_scale,
"per_act_token": per_act_token,
+ "ab_strides1": ab_strides1,
+ "ab_strides2": ab_strides2,
+ "c_strides1": c_strides1,
+ "c_strides2": c_strides2,
# cuda graph params
"cutlass_graph": cutlass_graph,
"triton_graph": triton_graph,
@@ -289,6 +318,10 @@ def bench_run(
w2_q,
w1_scale,
w2_scale,
+ ab_strides1,
+ ab_strides2,
+ c_strides1,
+ c_strides2,
topk_weights,
topk_ids,
per_act_token,
@@ -297,7 +330,7 @@ def bench_run(
results.append(
benchmark.Timer(
- stmt="run_cutlass_moe(a, a_scale, w1_q, w2_q, w1_scale, w2_scale, topk_weights, topk_ids, per_act_token, num_runs)", # noqa: E501
+ stmt="run_cutlass_moe(a, a_scale, w1_q, w2_q, w1_scale, w2_scale, ab_strides1, ab_strides2, c_strides1, c_strides2, topk_weights, topk_ids, per_act_token, num_runs)", # noqa: E501
globals=globals,
label=label,
sub_label=sub_label,
diff --git a/benchmarks/kernels/benchmark_lora.py b/benchmarks/kernels/benchmark_lora.py
index 3d38d4b353..89309c79f0 100644
--- a/benchmarks/kernels/benchmark_lora.py
+++ b/benchmarks/kernels/benchmark_lora.py
@@ -637,7 +637,7 @@ def bench_optype(
# Clear LoRA optimization hash-maps.
_LORA_A_PTR_DICT.clear()
_LORA_B_PTR_DICT.clear()
- # Run bench function so that _LORA_A_PTR_DICT and _LORA_B_PTR_DICT are setup
+ # Run bench function so that _LORA_A_PTR_DICT and _LORA_B_PTR_DICT are set up
for kwargs in kwargs_list:
op_type.bench_fn()(**kwargs)
torch.cuda.synchronize()
diff --git a/benchmarks/kernels/benchmark_machete.py b/benchmarks/kernels/benchmark_machete.py
index f73d0511e0..1b1c3b321c 100644
--- a/benchmarks/kernels/benchmark_machete.py
+++ b/benchmarks/kernels/benchmark_machete.py
@@ -236,6 +236,7 @@ def marlin_create_bench_fn(bt: BenchmarkTensors) -> Callable:
a=bt.a,
c=None,
b_q_weight=w_q,
+ b_bias=None,
b_scales=w_s,
global_scale=None,
b_zeros=w_zp,
@@ -252,28 +253,7 @@ def marlin_create_bench_fn(bt: BenchmarkTensors) -> Callable:
else:
assert bt.a.dtype == torch.int8
assert bt.wtype == scalar_types.uint4b8
-
- if bt.w_ch_s is not None:
- s_ch = bt.w_ch_s.to(torch.float32)
- else:
- s_ch = torch.ones(bt.w_ref.shape[1], dtype=torch.float32, device=device)
-
- if bt.w_tok_s is not None:
- s_tok = bt.w_tok_s.to(torch.float32)
- else:
- s_tok = torch.ones(bt.a.shape[0], dtype=torch.float32, device=device)
-
- fn = lambda: ops.marlin_qqq_gemm(
- a=bt.a,
- b_q_weight=w_q,
- s_group=w_s,
- s_tok=s_tok,
- s_ch=s_ch,
- workspace=workspace.scratch,
- size_m=bt.a.shape[0],
- size_n=bt.w_ref.shape[1],
- size_k=bt.w_ref.shape[0],
- )
+ raise NotImplementedError("QQQ is not supported anymore")
return fn
@@ -304,6 +284,25 @@ def machete_create_bench_fn(
)
+def cutlass_w4a8_create_bench_fn(
+ bt: BenchmarkTensors, out_type=torch.dtype, schedule=None
+) -> Callable:
+ w_q = bt.w_q.t().contiguous().t() # make col major
+ w_q = ops.cutlass_encode_and_reorder_int4b(w_q)
+ # expects fp8 scales
+ w_s = ops.cutlass_pack_scale_fp8(bt.w_g_s.to(torch.float8_e4m3fn))
+
+ return lambda: ops.cutlass_w4a8_mm(
+ a=bt.a,
+ b_q=w_q,
+ b_group_scales=w_s,
+ b_group_size=bt.group_size,
+ b_channel_scales=bt.w_ch_s,
+ a_token_scales=bt.w_tok_s,
+ maybe_schedule=schedule,
+ )
+
+
# impl
# bench
@@ -405,6 +404,20 @@ def bench(
)
)
+ # cutlass w4a8
+ if types.act_type == torch.float8_e4m3fn and group_size == 128:
+ timers.append(
+ bench_fns(
+ label,
+ sub_label,
+ f"cutlass w4a8 ({name_type_string})",
+ [
+ cutlass_w4a8_create_bench_fn(bt, out_type=types.output_type)
+ for bt in benchmark_tensors
+ ],
+ )
+ )
+
if sweep_schedules:
global _SWEEP_SCHEDULES_RESULTS
diff --git a/benchmarks/kernels/benchmark_moe.py b/benchmarks/kernels/benchmark_moe.py
index 72250e2fb6..6259aa0dd6 100644
--- a/benchmarks/kernels/benchmark_moe.py
+++ b/benchmarks/kernels/benchmark_moe.py
@@ -3,6 +3,7 @@
import argparse
import json
+import os
import time
from contextlib import nullcontext
from datetime import datetime
@@ -22,10 +23,10 @@ from vllm.utils import FlexibleArgumentParser
FP8_DTYPE = current_platform.fp8_dtype()
-def ensure_divisibility(numerator, denominator):
+def ensure_divisibility(numerator, denominator, text):
"""Ensure that numerator is divisible by the denominator."""
- assert numerator % denominator == 0, (
- "intermediate_size {} is not divisible by tp {}.".format(numerator, denominator)
+ assert numerator % denominator == 0, "{} {} is not divisible by tp {}.".format(
+ text, numerator, denominator
)
@@ -418,8 +419,10 @@ class BenchmarkWorker:
)
# NOTE(woosuk): The current naming convention uses w2.shape[2], which
# is the intermediate size after silu_and_mul.
+ block_n = block_quant_shape[0] if block_quant_shape else None
+ block_k = block_quant_shape[1] if block_quant_shape else None
op_config = get_moe_configs(
- num_experts, shard_intermediate_size // 2, dtype_str
+ num_experts, shard_intermediate_size // 2, dtype_str, block_n, block_k
)
if op_config is None:
config = get_default_config(
@@ -429,7 +432,7 @@ class BenchmarkWorker:
hidden_size,
topk,
dtype_str,
- is_marlin=False,
+ block_quant_shape,
)
else:
config = op_config[min(op_config.keys(), key=lambda x: abs(x - num_tokens))]
@@ -542,6 +545,7 @@ def save_configs(
use_fp8_w8a8: bool,
use_int8_w8a16: bool,
block_quant_shape: list[int],
+ save_dir: str,
) -> None:
dtype_str = get_config_dtype_str(
dtype, use_int8_w8a16=use_int8_w8a16, use_fp8_w8a8=use_fp8_w8a8
@@ -552,7 +556,8 @@ def save_configs(
filename = get_config_file_name(
num_experts, shard_intermediate_size // 2, dtype_str, block_quant_shape
)
-
+ os.makedirs(save_dir, exist_ok=True)
+ filename = os.path.join(save_dir, filename)
print(f"Writing best config to {filename}...")
with open(filename, "w") as f:
json.dump(configs, f, indent=4)
@@ -577,12 +582,10 @@ def main(args: argparse.Namespace):
E = config.ffn_config.moe_num_experts
topk = config.ffn_config.moe_top_k
intermediate_size = config.ffn_config.ffn_hidden_size
- shard_intermediate_size = 2 * intermediate_size // args.tp_size
elif config.architectures[0] == "JambaForCausalLM":
E = config.num_experts
topk = config.num_experts_per_tok
intermediate_size = config.intermediate_size
- shard_intermediate_size = 2 * intermediate_size // args.tp_size
elif config.architectures[0] in (
"DeepseekV3ForCausalLM",
"DeepseekV2ForCausalLM",
@@ -591,17 +594,14 @@ def main(args: argparse.Namespace):
E = config.n_routed_experts
topk = config.num_experts_per_tok
intermediate_size = config.moe_intermediate_size
- shard_intermediate_size = 2 * intermediate_size // args.tp_size
elif config.architectures[0] in ("Qwen2MoeForCausalLM", "Qwen3MoeForCausalLM"):
E = config.num_experts
topk = config.num_experts_per_tok
intermediate_size = config.moe_intermediate_size
- shard_intermediate_size = 2 * intermediate_size // args.tp_size
elif config.architectures[0] in ("HunYuanMoEV1ForCausalLM"):
E = config.num_experts
topk = config.moe_topk[0]
intermediate_size = config.moe_intermediate_size[0]
- shard_intermediate_size = 2 * intermediate_size // args.tp_size
else:
# Support for llama4
config = config.get_text_config()
@@ -609,8 +609,14 @@ def main(args: argparse.Namespace):
E = config.num_local_experts
topk = config.num_experts_per_tok
intermediate_size = config.intermediate_size
+ enable_ep = bool(args.enable_expert_parallel)
+ if enable_ep:
+ ensure_divisibility(E, args.tp_size, "Number of experts")
+ E = E // args.tp_size
+ shard_intermediate_size = 2 * intermediate_size
+ else:
+ ensure_divisibility(intermediate_size, args.tp_size, "intermediate_size")
shard_intermediate_size = 2 * intermediate_size // args.tp_size
- ensure_divisibility(intermediate_size, args.tp_size)
hidden_size = config.hidden_size
dtype = torch.float16 if current_platform.is_rocm() else config.torch_dtype
use_fp8_w8a8 = args.dtype == "fp8_w8a8"
@@ -672,7 +678,11 @@ def main(args: argparse.Namespace):
is_fp16 = not (use_fp8_w8a8 or use_int8_w8a16)
search_space = get_configs_compute_bound(is_fp16, block_quant_shape)
print(f"Start tuning over {len(search_space)} configurations...")
-
+ if use_deep_gemm:
+ raise ValueError(
+ "Tuning with --use-deep-gemm is not supported as it only tunes Triton "
+ "kernels. Please remove the flag."
+ )
start = time.time()
configs = _distribute(
"tune",
@@ -706,6 +716,7 @@ def main(args: argparse.Namespace):
use_fp8_w8a8,
use_int8_w8a16,
block_quant_shape,
+ args.save_dir,
)
end = time.time()
print(f"Tuning took {end - start:.2f} seconds")
@@ -742,10 +753,14 @@ if __name__ == "__main__":
parser.add_argument(
"--tp-size", "-tp", "--tensor-parallel-size", type=int, default=2
)
+ parser.add_argument("--enable-expert-parallel", "-enable-ep", action="store_true")
parser.add_argument(
"--dtype", type=str, choices=["auto", "fp8_w8a8", "int8_w8a16"], default="auto"
)
parser.add_argument("--use-deep-gemm", action="store_true")
+ parser.add_argument(
+ "--save-dir", type=str, default="./", help="Directory to save tuned results"
+ )
parser.add_argument("--seed", type=int, default=0)
parser.add_argument("--batch-size", type=int, nargs="+", required=False)
parser.add_argument("--tune", action="store_true")
diff --git a/benchmarks/kernels/benchmark_mrope.py b/benchmarks/kernels/benchmark_mrope.py
new file mode 100644
index 0000000000..b914736170
--- /dev/null
+++ b/benchmarks/kernels/benchmark_mrope.py
@@ -0,0 +1,328 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+# This script benchmarks the mrope kernel (mainly for Qwen2VL and Qwen2.5VL models).
+# It generates test data, runs benchmarks, and saves results to a CSV file.
+#
+# The CSV file (named with current date/time) contains these columns:
+# model_name, tp_size, num_tokens, num_heads, num_kv_heads, head_dim, max_position,
+# rope_theta, is_neox_style, rope_scaling, dtype, torch_mean, torch_median, torch_p99,
+# torch_min, torch_max, triton_mean, triton_median, triton_p99, triton_min, triton_max,
+# speedup
+#
+# == Usage Examples ==
+#
+# Single model benchmark:
+# python3 benchmark_mrope.py --model-name Qwen/Qwen2-VL-7B-Instruct --tp-size 1 \
+# --warmup-iter 10 --benchmark-iter 100 --dtype bfloat16 --seed 0 --num-tokens 1024
+#
+# All models benchmark:
+# python3 benchmark_mrope.py --model-name "" --tp-size 1 --warmup-iter 10 \
+# --benchmark-iter 100 --dtype bfloat16 --seed 0 --num-tokens 1024
+#
+# All models with different TP sizes:
+# python3 benchmark_mrope.py --model-name "" --tp-size 1 2 4 8 --warmup-iter 10 \
+# --benchmark-iter 100 --dtype bfloat16 --seed 0 --num-tokens 1024
+#
+# All models with different token counts:
+# python3 benchmark_mrope.py --model-name "" --tp-size 1 --warmup-iter 10 \
+# --benchmark-iter 100 --dtype bfloat16 --seed 0 --num-tokens 1024 4096 16384
+import csv
+import os
+import time
+from datetime import datetime
+from typing import Any
+
+import numpy as np
+import torch
+
+from vllm.model_executor.layers.rotary_embedding import get_rope
+from vllm.platforms import current_platform
+from vllm.transformers_utils.config import get_config
+from vllm.utils import FlexibleArgumentParser
+
+device = torch.device("cuda" if torch.cuda.is_available() else "cpu")
+
+
+def generate_test_data(
+ num_tokens: int,
+ num_q_heads: int,
+ num_kv_heads: int,
+ head_size: int,
+ max_position_embeddings: int,
+ dtype: torch.dtype,
+ device: torch.device,
+):
+ """Generate test data for given configuration."""
+ # Create 2D positions (3, num_tokens) for multimodal case
+ positions = torch.randint(
+ 0, max_position_embeddings // 4, (3, num_tokens), device=device
+ )
+
+ # Create query and key tensors
+ query = torch.randn(num_tokens, num_q_heads * head_size, dtype=dtype, device=device)
+ key = torch.randn(num_tokens, num_kv_heads * head_size, dtype=dtype, device=device)
+
+ return positions, query, key
+
+
+def calculate_stats(times: list[float]) -> dict[str, float]:
+ """Calculate statistics from a list of times."""
+ times_array = np.array(times)
+ return {
+ "mean": np.mean(times_array),
+ "median": np.median(times_array),
+ "p99": np.percentile(times_array, 99),
+ "min": np.min(times_array),
+ "max": np.max(times_array),
+ }
+
+
+def benchmark_mrope(
+ model_name: str,
+ num_tokens: int,
+ head_dim: int,
+ tp_size: int,
+ num_heads: int,
+ num_kv_heads: int,
+ max_position: int = 8192,
+ rope_theta: float = 10000,
+ is_neox_style: bool = True,
+ rope_scaling: dict[str, Any] = None,
+ dtype: torch.dtype = torch.bfloat16,
+ seed: int = 0,
+ warmup_iter: int = 10,
+ benchmark_iter: int = 100,
+ csv_writer=None,
+):
+ current_platform.seed_everything(seed)
+ torch.set_default_device(device)
+ # the parameters to compute the q k v size based on tp_size
+ mrope_helper_class = get_rope(
+ head_size=head_dim,
+ rotary_dim=head_dim,
+ max_position=max_position,
+ base=rope_theta,
+ is_neox_style=is_neox_style,
+ rope_scaling=rope_scaling,
+ dtype=dtype,
+ ).to(device=device)
+
+ print(80 * "=")
+ print(
+ f"Evaluating model: {model_name} "
+ f"with tp_size: {tp_size} "
+ f"and num_tokens: {num_tokens}, "
+ f"dtype: {dtype}"
+ )
+
+ # create q k v input tensors
+ # create rotary pos emb input tensors
+ positions, query, key = generate_test_data(
+ num_tokens, num_heads, num_kv_heads, head_dim, max_position, dtype, device
+ )
+
+ # Warm up
+ for _ in range(warmup_iter):
+ mrope_helper_class.forward_native(
+ positions,
+ query.clone(),
+ key.clone(),
+ )
+
+ mrope_helper_class.forward_cuda(
+ positions,
+ query.clone(),
+ key.clone(),
+ )
+
+ torch.cuda.synchronize()
+
+ # Time reference implementation
+ torch_times = []
+ for _ in range(benchmark_iter):
+ query_clone = query.clone()
+ key_clone = key.clone()
+ torch.cuda.synchronize()
+ start_time = time.time()
+
+ mrope_helper_class.forward_native(
+ positions,
+ query_clone,
+ key_clone,
+ )
+
+ torch.cuda.synchronize()
+ torch_times.append(time.time() - start_time)
+
+ # Time triton kernel implementation
+ triton_times = []
+ for _ in range(benchmark_iter):
+ query_clone = query.clone()
+ key_clone = key.clone()
+ torch.cuda.synchronize()
+ start_time = time.time()
+ mrope_helper_class.forward_cuda(
+ positions,
+ query_clone,
+ key_clone,
+ )
+ torch.cuda.synchronize()
+ triton_times.append(time.time() - start_time)
+
+ # Calculate statistics
+ torch_stats = calculate_stats(torch_times)
+ triton_stats = calculate_stats(triton_times)
+ print(f"\nPerformance for config ({num_tokens}, {num_heads}, {num_kv_heads}):")
+
+ print(
+ f"Torch implementation: "
+ f"mean={torch_stats['mean']:.8f}s, "
+ f"median={torch_stats['median']:.8f}s, "
+ f"p99={torch_stats['p99']:.8f}s"
+ )
+
+ print(
+ f"Triton implementation: "
+ f"mean={triton_stats['mean']:.8f}s, "
+ f"median={triton_stats['median']:.8f}s, "
+ f"p99={triton_stats['p99']:.8f}s"
+ )
+
+ print(
+ f"Triton Speedup over Torch: {torch_stats['mean'] / triton_stats['mean']:.8f}x"
+ )
+
+ # Write to CSV
+ if csv_writer:
+ row = [
+ model_name,
+ tp_size,
+ num_tokens,
+ num_heads,
+ num_kv_heads,
+ head_dim,
+ max_position,
+ rope_theta,
+ is_neox_style,
+ str(rope_scaling),
+ str(dtype).split(".")[-1],
+ torch_stats["mean"],
+ torch_stats["median"],
+ torch_stats["p99"],
+ torch_stats["min"],
+ torch_stats["max"],
+ triton_stats["mean"],
+ triton_stats["median"],
+ triton_stats["p99"],
+ triton_stats["min"],
+ triton_stats["max"],
+ torch_stats["mean"] / triton_stats["mean"], # speedup
+ ]
+ csv_writer.writerow(row)
+
+ return torch_stats, triton_stats
+
+
+if __name__ == "__main__":
+ parser = FlexibleArgumentParser(
+ description="Benchmark the rotary embedding kernels."
+ )
+ parser.add_argument("--model-name", type=str, default="")
+ parser.add_argument("--tp-size", type=int, default=1)
+ parser.add_argument("--warmup-iter", type=int, default=10)
+ parser.add_argument("--benchmark-iter", type=int, default=100)
+ parser.add_argument("--dtype", type=str, choices=["bfloat16"], default="bfloat16")
+ parser.add_argument("--seed", type=int, default=0)
+ parser.add_argument("--num-tokens", type=int, nargs="+", required=False)
+ parser.add_argument("--trust-remote-code", action="store_true")
+ parser.add_argument("--output-csv", type=str, default="mrope_benchmark_results.csv")
+ args = parser.parse_args()
+ print(args)
+
+ # Create CSV file for results
+ timestamp = datetime.now().strftime("%Y%m%d_%H%M%S")
+ csv_filename = f"{os.path.splitext(args.output_csv)[0]}_{timestamp}.csv"
+
+ with open(csv_filename, "w", newline="") as csvfile:
+ csv_writer = csv.writer(csvfile)
+ # Write header
+ header = [
+ "model_name",
+ "tp_size",
+ "num_tokens",
+ "num_heads",
+ "num_kv_heads",
+ "head_dim",
+ "max_position",
+ "rope_theta",
+ "is_neox_style",
+ "rope_scaling",
+ "dtype",
+ "torch_mean",
+ "torch_median",
+ "torch_p99",
+ "torch_min",
+ "torch_max",
+ "triton_mean",
+ "triton_median",
+ "triton_p99",
+ "triton_min",
+ "triton_max",
+ "speedup",
+ ]
+ csv_writer.writerow(header)
+
+ model_tp_dict = {}
+ if args.model_name == "":
+ model_tp_dict = {
+ "Qwen/Qwen2-VL-2B-Instruct": [1],
+ "Qwen/Qwen2-VL-7B-Instruct": [1],
+ "Qwen/Qwen2-VL-72B-Instruct": [2, 4, 8],
+ "Qwen/Qwen2.5-VL-3B-Instruct": [1, 2, 4, 8],
+ "Qwen/Qwen2.5-VL-7B-Instruct": [1, 2, 4, 8],
+ "Qwen/Qwen2.5-VL-72B-Instruct": [2, 4, 8],
+ }
+ else:
+ model_tp_dict[args.model_name] = [args.tp_size]
+
+ if args.num_tokens is None:
+ num_tokens_list = [2**i for i in range(0, 18)]
+ else:
+ num_tokens_list = args.num_tokens
+
+ for model_name, tp_list in model_tp_dict.items():
+ config = get_config(model_name, trust_remote_code=args.trust_remote_code)
+ for tp_size in tp_list:
+ # get the model config
+ total_num_kv_heads = config.num_key_value_heads
+ total_num_heads = config.num_attention_heads
+ num_heads = total_num_heads // tp_size
+ num_kv_heads = max(1, total_num_kv_heads // tp_size)
+ head_dim = config.hidden_size // total_num_heads
+ q_size = num_heads * head_dim
+ kv_size = num_kv_heads * head_dim
+ is_neox_style = True
+ rope_theta = config.rope_theta
+ max_position = config.max_position_embeddings
+
+ for num_tokens in num_tokens_list:
+ benchmark_mrope(
+ model_name=model_name,
+ num_tokens=num_tokens,
+ head_dim=head_dim,
+ tp_size=tp_size,
+ num_heads=num_heads,
+ num_kv_heads=num_kv_heads,
+ max_position=max_position,
+ rope_theta=rope_theta,
+ is_neox_style=is_neox_style,
+ rope_scaling=config.rope_scaling,
+ dtype=getattr(torch, args.dtype),
+ seed=args.seed,
+ warmup_iter=args.warmup_iter,
+ benchmark_iter=args.benchmark_iter,
+ csv_writer=csv_writer,
+ )
+
+ print(f"Benchmark results saved to {csv_filename}")
diff --git a/benchmarks/kernels/benchmark_silu_mul_fp8_quant.py b/benchmarks/kernels/benchmark_silu_mul_fp8_quant.py
new file mode 100644
index 0000000000..0650cbf3cc
--- /dev/null
+++ b/benchmarks/kernels/benchmark_silu_mul_fp8_quant.py
@@ -0,0 +1,77 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import time
+
+import torch
+
+from vllm.model_executor.layers.fused_moe.batched_deep_gemm_moe import (
+ silu_mul_fp8_quant_deep_gemm,
+)
+from vllm.platforms import current_platform
+
+
+def benchmark(E, T, H, G=128, runs=50):
+ current_platform.seed_everything(42)
+ y = torch.randn((E, T, 2 * H), dtype=torch.bfloat16, device="cuda")
+ tokens_per_expert = torch.randint(
+ T // 2, T, size=(E,), dtype=torch.int32, device="cuda"
+ )
+
+ # Warmup
+ for _ in range(10):
+ silu_mul_fp8_quant_deep_gemm(y, tokens_per_expert, group_size=G)
+ torch.cuda.synchronize()
+
+ # Benchmark
+ torch.cuda.synchronize()
+ start = time.perf_counter()
+ for _ in range(runs):
+ silu_mul_fp8_quant_deep_gemm(y, tokens_per_expert, group_size=G)
+ torch.cuda.synchronize()
+
+ avg_time = (time.perf_counter() - start) / runs * 1000
+
+ # Calculate actual work done (only count valid tokens)
+ actual_tokens = tokens_per_expert.sum().item()
+ actual_elements = actual_tokens * H
+
+ # GFLOPS: operations per element = exp + 3 muls + 1 div + quantization ops ≈ 8 ops
+ ops_per_element = 8
+ total_ops = actual_elements * ops_per_element
+ gflops = total_ops / (avg_time / 1000) / 1e9
+
+ # Memory bandwidth: bfloat16 inputs (2 bytes), fp8 output (1 byte), scales (4 bytes)
+ input_bytes = actual_tokens * 2 * H * 2 # 2*H bfloat16 inputs
+ output_bytes = actual_tokens * H * 1 # H fp8 outputs
+ scale_bytes = actual_tokens * (H // G) * 4 # scales in float32
+ total_bytes = input_bytes + output_bytes + scale_bytes
+ memory_bw = total_bytes / (avg_time / 1000) / 1e9
+
+ return avg_time, gflops, memory_bw
+
+
+configs = [
+ (8, 32, 1024),
+ (16, 64, 2048),
+ (32, 128, 4096),
+ # DeepSeekV3 Configs
+ (256, 16, 7168),
+ (256, 32, 7168),
+ (256, 64, 7168),
+ (256, 128, 7168),
+ (256, 256, 7168),
+ (256, 512, 7168),
+ (256, 1024, 7168),
+]
+
+print(f"GPU: {torch.cuda.get_device_name()}")
+print(f"{'Config':<20} {'Time(ms)':<10} {'GFLOPS':<10} {'GB/s':<10}")
+print("-" * 50)
+
+for E, T, H in configs:
+ try:
+ time_ms, gflops, gbps = benchmark(E, T, H)
+ print(f"E={E:3d},T={T:4d},H={H:4d} {time_ms:8.3f} {gflops:8.1f} {gbps:8.1f}")
+ except Exception:
+ print(f"E={E:3d},T={T:4d},H={H:4d} FAILED")
diff --git a/benchmarks/kernels/benchmark_trtllm_decode_attention.py b/benchmarks/kernels/benchmark_trtllm_decode_attention.py
index 77136edca4..603ce5ecf0 100644
--- a/benchmarks/kernels/benchmark_trtllm_decode_attention.py
+++ b/benchmarks/kernels/benchmark_trtllm_decode_attention.py
@@ -3,16 +3,17 @@
import csv
import os
-import random
from datetime import datetime
+from typing import Optional
import flashinfer
import torch
-FLOAT32_BYTES = torch.finfo(torch.float).bits // 8
+from vllm.utils import round_up
-# KV Cache Layout for TRT-LLM
-# kv_cache_shape = (num_blocks, 2, num_kv_heads, page_size, head_dim)
+FLOAT32_BYTES = torch.finfo(torch.float).bits // 8
+FP8_DTYPE = torch.float8_e4m3fn
+FP4_DTYPE = torch.uint8
def to_float8(x, dtype=torch.float8_e4m3fn):
@@ -26,65 +27,106 @@ def to_float8(x, dtype=torch.float8_e4m3fn):
@torch.no_grad()
def benchmark_decode(
- num_seqs,
- max_seq_len,
- page_size=16,
- dtype=torch.bfloat16,
- kv_layout="HND",
- num_kv_heads=8,
- kv_cache_dtype="auto",
- head_dim=128,
- warmup=10,
- trials=20,
+ dtype: torch.dtype,
+ quant_dtypes: tuple[
+ Optional[torch.dtype], Optional[torch.dtype], Optional[torch.dtype]
+ ],
+ batch_size: int,
+ max_seq_len: int,
+ num_heads: tuple[int, int] = (64, 8),
+ head_size: int = 128,
+ kv_layout: str = "HND",
+ block_size: int = 16,
+ warmup: int = 10,
+ trials: int = 20,
):
torch.set_default_device("cuda")
- device = "cuda"
torch.manual_seed(0)
- HEAD_GRP_SIZE = 8
- MAX_SEQ_LEN = max_seq_len
+ q_quant_dtype, kv_quant_dtype, o_quant_dtype = quant_dtypes
+ q_quant_dtype = q_quant_dtype or dtype
+ kv_quant_dtype = kv_quant_dtype or dtype
+ o_quant_dtype = o_quant_dtype or dtype
+
+ num_qo_heads, num_kv_heads = num_heads
+ assert num_qo_heads % num_kv_heads == 0
+
+ sm_scale = float(1.0 / (head_size**0.5))
# large number to reduce kv_cache reuse
- NUM_BLOCKS = int(256000 / page_size)
+ NUM_BLOCKS = int(256000 / block_size)
- workspace_buffer = torch.empty(1024 * 1024 * 1024, dtype=torch.int8, device=device)
+ kv_cache_shape = None
+ if kv_layout == "NHD":
+ kv_cache_shape = (NUM_BLOCKS, 2, block_size, num_kv_heads, head_size)
+ elif kv_layout == "HND":
+ kv_cache_shape = (NUM_BLOCKS, 2, num_kv_heads, block_size, head_size)
+ else:
+ raise ValueError(f"Invalid kv_layout: {kv_layout}")
- # For decode, batch_size is num_decode_token
- num_qo_heads = num_kv_heads * HEAD_GRP_SIZE
- sm_scale = float(1.0 / (head_dim**0.5))
- q = torch.randn(num_seqs, num_qo_heads, head_dim, device=device, dtype=dtype)
- kv_lens = [random.randint(1, MAX_SEQ_LEN) for _ in range(num_seqs)]
+ # Always using 1.0 scale to reflect the real perf in benchmarking
+ q_scale = 1.0
+ ref_query = torch.randn(batch_size, num_qo_heads, head_size, dtype=dtype)
+ if q_quant_dtype == FP8_DTYPE:
+ query, _ = to_float8(ref_query)
+ else:
+ query = ref_query
- max_kv_len = max(kv_lens)
- kv_lens_tensor = torch.tensor(kv_lens, dtype=torch.int, device=device)
- max_num_blocks_per_seq = (max_kv_len + page_size - 1) // page_size
+ kv_lens = torch.randint(1, max_seq_len, (batch_size,), dtype=torch.int32)
+ kv_lens[-1] = max_seq_len
- block_tables = torch.randint(
- 0, NUM_BLOCKS, (num_seqs, max_num_blocks_per_seq), dtype=torch.int32
- )
+ seq_lens = kv_lens
+ max_seq_len = torch.max(seq_lens).item()
- kv_cache_shape = (NUM_BLOCKS, 2, num_kv_heads, page_size, head_dim)
- kv_cache = torch.randn(size=kv_cache_shape, device=device, dtype=dtype)
+ # Always using 1.0 scale to reflect the real perf in benchmarking
k_scale = v_scale = 1.0
+ ref_kv_cache = torch.randn(kv_cache_shape, dtype=dtype)
+ if kv_quant_dtype == FP8_DTYPE:
+ kv_cache, _ = to_float8(ref_kv_cache)
+ else:
+ kv_cache = ref_kv_cache
- if kv_cache_dtype.startswith("fp8"):
- kv_cache, _ = to_float8(kv_cache)
+ max_num_blocks_per_seq = (max_seq_len + block_size - 1) // block_size
+ block_tables = torch.randint(
+ 0, NUM_BLOCKS, (batch_size, max_num_blocks_per_seq), dtype=torch.int32
+ )
+ kv_indptr = [0]
+ kv_indices = []
+ kv_last_page_lens = []
+ for i in range(batch_size):
+ seq_len = seq_lens[i]
+ assert seq_len > 0
+ num_blocks = (seq_len + block_size - 1) // block_size
+ kv_indices.extend(block_tables[i, :num_blocks])
+ kv_indptr.append(kv_indptr[-1] + num_blocks)
+ kv_last_page_len = seq_len % block_size
+ if kv_last_page_len == 0:
+ kv_last_page_len = block_size
+ kv_last_page_lens.append(kv_last_page_len)
- output_trtllm = torch.empty(q.shape, dtype=dtype)
+ kv_indptr = torch.tensor(kv_indptr, dtype=torch.int32)
+ kv_indices = torch.tensor(kv_indices, dtype=torch.int32)
+ kv_last_page_lens = torch.tensor(kv_last_page_lens, dtype=torch.int32)
+ workspace_buffer = torch.zeros(1024 * 1024 * 1024, dtype=torch.int8)
- # Benchmark TRT decode
- def trt_decode():
- return flashinfer.decode.trtllm_batch_decode_with_kv_cache(
- q,
- kv_cache,
- workspace_buffer,
- block_tables,
- kv_lens_tensor,
- max_kv_len,
- bmm1_scale=k_scale * sm_scale,
- bmm2_scale=v_scale,
- out=output_trtllm,
- )
+ wrapper = flashinfer.BatchDecodeWithPagedKVCacheWrapper(
+ workspace_buffer,
+ kv_layout,
+ use_tensor_cores=True,
+ )
+ wrapper.plan(
+ kv_indptr,
+ kv_indices,
+ kv_last_page_lens,
+ num_qo_heads,
+ num_kv_heads,
+ head_size,
+ block_size,
+ "NONE",
+ sm_scale=sm_scale,
+ q_data_type=dtype,
+ kv_data_type=dtype,
+ )
def time_fn(fn, warmup=10, trials=20):
torch.cuda.synchronize()
@@ -101,74 +143,72 @@ def benchmark_decode(
times.append(start.elapsed_time(end)) # ms
return sum(times) / len(times), torch.std(torch.tensor(times))
- # TRT Decode
- trt_mean, trt_std = time_fn(trt_decode)
-
- kv_indptr = [0]
- kv_indices = []
- kv_last_page_lens = []
- for i in range(num_seqs):
- seq_len = kv_lens[i]
- assert seq_len > 0
- num_blocks = (seq_len + page_size - 1) // page_size
- kv_indices.extend(block_tables[i, :num_blocks])
- kv_indptr.append(kv_indptr[-1] + num_blocks)
- kv_last_page_len = seq_len % page_size
- if kv_last_page_len == 0:
- kv_last_page_len = page_size
- kv_last_page_lens.append(kv_last_page_len)
-
- kv_indptr = torch.tensor(kv_indptr, dtype=torch.int32)
- kv_indices = torch.tensor(kv_indices, dtype=torch.int32)
- kv_last_page_lens = torch.tensor(kv_last_page_lens, dtype=torch.int32)
-
- output_baseline = torch.empty(q.shape, dtype=dtype)
-
- wrapper = flashinfer.BatchDecodeWithPagedKVCacheWrapper(
- workspace_buffer,
- kv_layout,
- use_tensor_cores=((num_qo_heads // num_kv_heads) > 4),
- )
-
- wrapper.plan(
- kv_indptr,
- kv_indices,
- kv_last_page_lens,
- num_qo_heads,
- num_kv_heads,
- head_dim,
- page_size,
- "NONE",
- q_data_type=dtype,
- kv_data_type=torch.float8_e4m3fn if kv_cache_dtype.startswith("fp8") else dtype,
- )
+ o_scale = 1.0
+ o_sf_scale = None
+ output_baseline = torch.empty(ref_query.shape, dtype=dtype)
+ if o_quant_dtype == FP4_DTYPE:
+ o_sf_scale = 500.0
+ output_trtllm = flashinfer.utils.FP4Tensor(
+ torch.empty(query.shape[:-1] + (query.shape[-1] // 2,), dtype=torch.uint8),
+ torch.empty(
+ (
+ round_up(query.shape[0], 128),
+ round_up(query.shape[1] * query.shape[2] // 16, 4),
+ ),
+ dtype=torch.float8_e4m3fn,
+ ),
+ )
+ else:
+ output_trtllm = torch.empty(query.shape, dtype=o_quant_dtype)
def baseline_decode():
- return wrapper.run(q, kv_cache, sm_scale, k_scale, v_scale, output_baseline)
+ return wrapper.run(
+ ref_query,
+ ref_kv_cache,
+ k_scale=k_scale,
+ v_scale=v_scale,
+ out=output_baseline,
+ )
+
+ def trtllm_decode():
+ return flashinfer.decode.trtllm_batch_decode_with_kv_cache(
+ query=query,
+ kv_cache=kv_cache,
+ workspace_buffer=workspace_buffer,
+ block_tables=block_tables,
+ seq_lens=seq_lens,
+ max_seq_len=max_seq_len,
+ bmm1_scale=q_scale * k_scale * sm_scale,
+ bmm2_scale=v_scale / o_scale,
+ o_sf_scale=o_sf_scale,
+ out=output_trtllm,
+ )
baseline_mean, baseline_std = time_fn(baseline_decode)
+ trtllm_mean, trtllm_std = time_fn(trtllm_decode)
# Calculate percentage speedup (positive means TRT is faster)
- speedup_percent = (baseline_mean - trt_mean) / baseline_mean
+ speedup_percent = (baseline_mean - trtllm_mean) / baseline_mean
print(
- f"\t{num_seqs}\t{max_seq_len}\t{trt_mean:.3f}\t{trt_std.item():.3f}"
+ f"\t{batch_size}\t{max_seq_len}\t{trtllm_mean:.3f}\t{trtllm_std.item():.3f}"
f"\t{baseline_mean:.3f}\t{baseline_std.item():.3f}\t{speedup_percent:.3f}"
)
# Return results for CSV writing
return {
- "num_seqs": num_seqs,
- "trt_mean": trt_mean,
- "trt_std": trt_std.item(),
+ "batch_size": batch_size,
+ "trtllm_mean": trtllm_mean,
+ "trtllm_std": trtllm_std.item(),
"baseline_mean": baseline_mean,
"baseline_std": baseline_std.item(),
"speedup_percent": speedup_percent,
- "q_dtype": str(dtype),
- "kv_cache_dtype": kv_cache_dtype,
- "page_size": page_size,
+ "q_dtype": str(q_quant_dtype),
+ "kv_cache_dtype": str(kv_quant_dtype),
+ "output_dtype": str(o_quant_dtype),
+ "block_size": block_size,
"num_kv_heads": num_kv_heads,
- "head_dim": head_dim,
+ "head_size": head_size,
"max_seq_len": max_seq_len,
}
@@ -180,17 +220,18 @@ def write_results_to_csv(results, filename=None):
filename = f"flashinfer_trtllm_benchmark_{timestamp}.csv"
fieldnames = [
- "num_seqs",
- "trt_mean",
- "trt_std",
+ "batch_size",
+ "trtllm_mean",
+ "trtllm_std",
"baseline_mean",
"baseline_std",
"speedup_percent",
"q_dtype",
"kv_cache_dtype",
- "page_size",
+ "output_dtype",
+ "block_size",
"num_kv_heads",
- "head_dim",
+ "head_size",
"max_seq_len",
]
@@ -209,45 +250,43 @@ def write_results_to_csv(results, filename=None):
if __name__ == "__main__":
- num_seqs = [1, 4, 8, 16, 32, 64, 128, 256]
+ batch_sizes = [1, 4, 8, 16, 32, 64, 128, 256]
max_seq_lens = [1024, 2048, 4096, 8192, 16384, 32768, 65536, 131072]
all_results = []
- print(
- "Running benchmark for q_dtype = bfloat16, kv_cache_dtype: bfloat16, "
- "output_dtype: bfloat16"
- )
- print(
- "\tnum_seqs\tmax_seq_len\ttrt_mean\ttrt_std\tbaseline_mean\t"
- "baseline_std\tspeedup_percent"
- )
- for max_seq_len in max_seq_lens:
- for bs in num_seqs:
- result = benchmark_decode(
- bs,
- max_seq_len,
- dtype=torch.bfloat16,
- kv_cache_dtype="auto",
- )
- all_results.append(result)
+ dtype = torch.bfloat16
+ quant_dtypes = [
+ # (q_quant_dtype, kv_quant_dtype, o_quant_dtype)
+ (None, None, None),
+ (None, FP8_DTYPE, None),
+ (FP8_DTYPE, FP8_DTYPE, FP8_DTYPE),
+ (FP8_DTYPE, FP8_DTYPE, FP4_DTYPE),
+ ]
- print(
- "Running benchmark for q_dtype = bfloat16, kv_cache_dtype: fp8, "
- "output_dtype: bfloat16"
- )
- print(
- "\tnum_seqs\tmax_seq_len\ttrt_mean\ttrt_std\tbaseline_mean\t"
- "baseline_std\tspeedup_percent"
- )
- for max_seq_len in max_seq_lens:
- for bs in num_seqs:
- result = benchmark_decode(
- bs,
- max_seq_len,
- dtype=torch.bfloat16,
- kv_cache_dtype="fp8",
- )
- all_results.append(result)
+ for quant_dtype in quant_dtypes:
+ q_quant_dtype, kv_quant_dtype, o_quant_dtype = quant_dtype
+ q_quant_dtype = q_quant_dtype or dtype
+ kv_quant_dtype = kv_quant_dtype or dtype
+ o_quant_dtype = o_quant_dtype or dtype
+
+ print(
+ f"Running benchmark for q_dtype = {q_quant_dtype}, "
+ f"kv_cache_dtype: {kv_quant_dtype}, "
+ f"output_dtype: {o_quant_dtype}"
+ )
+ print(
+ "\tbatch_size\tmax_seq_len\ttrtllm_mean\ttrtllm_std\tbaseline_mean\t"
+ "baseline_std\tspeedup_percent"
+ )
+ for max_seq_len in max_seq_lens:
+ for bs in batch_sizes:
+ result = benchmark_decode(
+ dtype=dtype,
+ quant_dtypes=quant_dtype,
+ batch_size=bs,
+ max_seq_len=max_seq_len,
+ )
+ all_results.append(result)
# Write all results to CSV
write_results_to_csv(all_results)
diff --git a/benchmarks/kernels/benchmark_trtllm_prefill_attention.py b/benchmarks/kernels/benchmark_trtllm_prefill_attention.py
index 67bd9aebbc..40903c6c34 100644
--- a/benchmarks/kernels/benchmark_trtllm_prefill_attention.py
+++ b/benchmarks/kernels/benchmark_trtllm_prefill_attention.py
@@ -3,16 +3,17 @@
import csv
import os
-import random
from datetime import datetime
+from typing import Optional
import flashinfer
import torch
-FLOAT32_BYTES = torch.finfo(torch.float).bits // 8
+from vllm.utils import round_up
-# KV Cache Layout for TRT-LLM
-# kv_cache_shape = (num_blocks, 2, num_kv_heads, page_size, head_dim)
+FLOAT32_BYTES = torch.finfo(torch.float).bits // 8
+FP8_DTYPE = torch.float8_e4m3fn
+FP4_DTYPE = torch.uint8
def to_float8(x, dtype=torch.float8_e4m3fn):
@@ -26,84 +27,100 @@ def to_float8(x, dtype=torch.float8_e4m3fn):
@torch.no_grad()
def benchmark_prefill(
- num_seqs,
- max_seq_len,
- page_size=16,
- dtype=torch.bfloat16,
- kv_layout="HND",
- num_kv_heads=8,
- kv_cache_dtype="auto",
- head_dim=128,
- warmup=10,
- trials=20,
+ dtype: torch.dtype,
+ quant_dtypes: tuple[
+ Optional[torch.dtype], Optional[torch.dtype], Optional[torch.dtype]
+ ],
+ batch_size: int,
+ max_seq_len: int,
+ num_heads: tuple[int, int] = (64, 8),
+ head_size: int = 128,
+ kv_layout: str = "HND",
+ block_size: int = 16,
+ warmup: int = 10,
+ trials: int = 20,
):
torch.set_default_device("cuda")
torch.manual_seed(0)
- HEAD_GRP_SIZE = 8
- MAX_SEQ_LEN = max_seq_len
+ q_quant_dtype, kv_quant_dtype, o_quant_dtype = quant_dtypes
+ q_quant_dtype = q_quant_dtype or dtype
+ kv_quant_dtype = kv_quant_dtype or dtype
+ o_quant_dtype = o_quant_dtype or dtype
+
+ max_q_len = max_kv_len = max_seq_len
+
+ num_qo_heads, num_kv_heads = num_heads
+ assert num_qo_heads % num_kv_heads == 0
+
+ sm_scale = float(1.0 / (head_size**0.5))
# large number to reduce kv_cache reuse
- NUM_BLOCKS = int(256000 / page_size)
+ NUM_BLOCKS = int(256000 / block_size)
- workspace_buffer = torch.empty(1024 * 1024 * 1024, dtype=torch.int8)
+ kv_cache_shape = None
+ if kv_layout == "NHD":
+ kv_cache_shape = (NUM_BLOCKS, 2, block_size, num_kv_heads, head_size)
+ elif kv_layout == "HND":
+ kv_cache_shape = (NUM_BLOCKS, 2, num_kv_heads, block_size, head_size)
+ else:
+ raise ValueError(f"Invalid kv_layout: {kv_layout}")
- num_qo_heads = num_kv_heads * HEAD_GRP_SIZE
- sm_scale = float(1.0 / (head_dim**0.5))
-
- q_lens = [random.randint(1, MAX_SEQ_LEN) for _ in range(num_seqs)]
- q_lens[-1] = MAX_SEQ_LEN
- max_q_len = max(q_lens)
+ q_lens = torch.randint(1, max_q_len, (batch_size,), dtype=torch.int32)
+ q_lens[-1] = max_q_len
q_indptr = torch.cat(
[
torch.tensor([0], dtype=torch.int32),
- torch.cumsum(
- torch.tensor(q_lens, dtype=torch.int32), dim=0, dtype=torch.int32
- ),
+ torch.cumsum(q_lens, dim=0, dtype=torch.int32),
]
)
- q = torch.randn(sum(q_lens), num_qo_heads, head_dim, dtype=dtype)
- kv_lens = [random.randint(0, MAX_SEQ_LEN) for _ in range(num_seqs)]
- kv_lens[-1] = MAX_SEQ_LEN
-
- seq_lens = [q_len + kv_len for q_len, kv_len in zip(q_lens, kv_lens)]
- max_seq_len = max(seq_lens)
- seq_lens_tensor = torch.tensor(seq_lens, dtype=torch.int32)
-
- max_num_blocks_per_seq = (max_seq_len + page_size - 1) // page_size
- block_tables = torch.randint(
- 0, NUM_BLOCKS, (num_seqs, max_num_blocks_per_seq), dtype=torch.int32
+ # Always using 1.0 scale to reflect the real perf in benchmarking
+ q_scale = 1.0
+ ref_query = torch.randn(
+ torch.sum(q_lens).item(), num_qo_heads, head_size, dtype=dtype
)
+ if q_quant_dtype == FP8_DTYPE:
+ query, _ = to_float8(ref_query)
+ else:
+ query = ref_query
- kv_cache_shape = (NUM_BLOCKS, 2, num_kv_heads, page_size, head_dim)
- kv_cache = torch.randn(size=kv_cache_shape, dtype=dtype)
+ kv_lens = torch.randint(0, max_kv_len, (batch_size,), dtype=torch.int32)
+ kv_lens[-1] = max_kv_len
+
+ seq_lens = kv_lens + q_lens
+ max_seq_len = torch.max(seq_lens).item()
+
+ # Always using 1.0 scale to reflect the real perf in benchmarking
k_scale = v_scale = 1.0
+ ref_kv_cache = torch.randn(kv_cache_shape, dtype=dtype)
+ if kv_quant_dtype == FP8_DTYPE:
+ kv_cache, _ = to_float8(ref_kv_cache)
+ else:
+ kv_cache = ref_kv_cache
- if kv_cache_dtype.startswith("fp8"):
- kv_cache, _ = to_float8(kv_cache)
-
- output_trtllm = torch.empty(q.shape, dtype=dtype)
-
+ max_num_blocks_per_seq = (max_seq_len + block_size - 1) // block_size
+ block_tables = torch.randint(
+ 0, NUM_BLOCKS, (batch_size, max_num_blocks_per_seq), dtype=torch.int32
+ )
kv_indptr = [0]
kv_indices = []
kv_last_page_lens = []
- for i in range(num_seqs):
+ for i in range(batch_size):
seq_len = seq_lens[i]
assert seq_len > 0
- num_blocks = (seq_len + page_size - 1) // page_size
+ num_blocks = (seq_len + block_size - 1) // block_size
kv_indices.extend(block_tables[i, :num_blocks])
kv_indptr.append(kv_indptr[-1] + num_blocks)
- kv_last_page_len = seq_len % page_size
+ kv_last_page_len = seq_len % block_size
if kv_last_page_len == 0:
- kv_last_page_len = page_size
+ kv_last_page_len = block_size
kv_last_page_lens.append(kv_last_page_len)
kv_indptr = torch.tensor(kv_indptr, dtype=torch.int32)
kv_indices = torch.tensor(kv_indices, dtype=torch.int32)
kv_last_page_lens = torch.tensor(kv_last_page_lens, dtype=torch.int32)
-
- output_baseline = torch.empty(q.shape, dtype=dtype)
+ workspace_buffer = torch.zeros(1024 * 1024 * 1024, dtype=torch.int8)
wrapper = flashinfer.BatchPrefillWithPagedKVCacheWrapper(
workspace_buffer, kv_layout
@@ -115,12 +132,12 @@ def benchmark_prefill(
kv_last_page_lens,
num_qo_heads,
num_kv_heads,
- head_dim,
- page_size,
+ head_size,
+ block_size,
causal=True,
sm_scale=sm_scale,
q_data_type=dtype,
- kv_data_type=kv_cache.dtype,
+ kv_data_type=dtype,
)
def time_fn(fn, warmup=10, trials=20):
@@ -138,52 +155,76 @@ def benchmark_prefill(
times.append(start.elapsed_time(end)) # ms
return sum(times) / len(times), torch.std(torch.tensor(times))
+ o_scale = 1.0
+ o_sf_scale = None
+ output_baseline = torch.empty(ref_query.shape, dtype=dtype)
+ if o_quant_dtype == FP4_DTYPE:
+ o_sf_scale = 500.0
+ output_trtllm = flashinfer.utils.FP4Tensor(
+ torch.empty(query.shape[:-1] + (query.shape[-1] // 2,), dtype=torch.uint8),
+ torch.empty(
+ (
+ round_up(query.shape[0], 128),
+ round_up(query.shape[1] * query.shape[2] // 16, 4),
+ ),
+ dtype=torch.float8_e4m3fn,
+ ),
+ )
+ else:
+ output_trtllm = torch.empty(query.shape, dtype=o_quant_dtype)
+
def baseline_prefill():
return wrapper.run(
- q, kv_cache, k_scale=k_scale, v_scale=v_scale, out=output_baseline
+ ref_query,
+ ref_kv_cache,
+ k_scale=k_scale,
+ v_scale=v_scale,
+ out=output_baseline,
)
- def trt_prefill():
+ def trtllm_prefill():
return flashinfer.prefill.trtllm_batch_context_with_kv_cache(
- query=q,
+ query=query,
kv_cache=kv_cache,
workspace_buffer=workspace_buffer,
block_tables=block_tables,
- seq_lens=seq_lens_tensor,
+ seq_lens=seq_lens,
max_q_len=max_q_len,
max_kv_len=max_seq_len,
- bmm1_scale=k_scale * sm_scale,
- bmm2_scale=v_scale,
- batch_size=num_seqs,
+ bmm1_scale=q_scale * k_scale * sm_scale,
+ bmm2_scale=v_scale / o_scale,
+ batch_size=batch_size,
cum_seq_lens_q=q_indptr,
cum_seq_lens_kv=kv_indptr,
+ o_sf_scale=o_sf_scale,
out=output_trtllm,
)
- trt_mean, trt_std = time_fn(trt_prefill)
baseline_mean, baseline_std = time_fn(baseline_prefill)
+ trtllm_mean, trtllm_std = time_fn(trtllm_prefill)
# Calculate percentage speedup (positive means TRT is faster)
- speedup_percent = (baseline_mean - trt_mean) / baseline_mean
+ speedup_percent = (baseline_mean - trtllm_mean) / baseline_mean
print(
- f"\t{num_seqs}\t{max_seq_len}\t{trt_mean:.5f}\t{trt_std.item():.5f}"
- f"\t{baseline_mean:.5f}\t{baseline_std.item():.5f}\t{speedup_percent:.5f}"
+ f"\t{batch_size}\t{max_seq_len}\t{trtllm_mean:8.3f}\t{trtllm_std.item():8.3f}"
+ f"\t{baseline_mean:8.3f}\t{baseline_std.item():8.3f}\t{speedup_percent:8.3f}"
)
# Return results for CSV writing
return {
- "num_seqs": num_seqs,
- "trt_mean": trt_mean,
- "trt_std": trt_std.item(),
+ "batch_size": batch_size,
+ "trtllm_mean": trtllm_mean,
+ "trtllm_std": trtllm_std.item(),
"baseline_mean": baseline_mean,
"baseline_std": baseline_std.item(),
"speedup_percent": speedup_percent,
- "q_dtype": str(dtype),
- "kv_cache_dtype": kv_cache_dtype,
- "page_size": page_size,
+ "q_dtype": str(q_quant_dtype),
+ "kv_cache_dtype": str(kv_quant_dtype),
+ "output_dtype": str(o_quant_dtype),
+ "block_size": block_size,
"num_kv_heads": num_kv_heads,
- "head_dim": head_dim,
+ "head_size": head_size,
"max_seq_len": max_seq_len,
}
@@ -195,17 +236,18 @@ def write_results_to_csv(results, filename=None):
filename = f"flashinfer_trtllm_benchmark_{timestamp}.csv"
fieldnames = [
- "num_seqs",
- "trt_mean",
- "trt_std",
+ "batch_size",
+ "trtllm_mean",
+ "trtllm_std",
"baseline_mean",
"baseline_std",
"speedup_percent",
"q_dtype",
"kv_cache_dtype",
- "page_size",
+ "output_dtype",
+ "block_size",
"num_kv_heads",
- "head_dim",
+ "head_size",
"max_seq_len",
]
@@ -224,27 +266,42 @@ def write_results_to_csv(results, filename=None):
if __name__ == "__main__":
- num_seqs = [1, 4, 8, 16, 32, 64, 128, 256]
+ batch_sizes = [1, 4, 8, 16, 32, 64, 128, 256]
max_seq_lens = [1024, 2048, 4096, 8192, 16384, 32768, 65536, 131072]
all_results = []
- print(
- "Running benchmark for q_dtype = bfloat16, kv_cache_dtype: bfloat16, "
- "output_dtype: bfloat16"
- )
- print(
- "\tnum_seqs\tmax_seq_len\ttrt_mean\ttrt_std\tbaseline_mean\t"
- "baseline_std\tspeedup_percent"
- )
- for max_seq_len in max_seq_lens:
- for bs in num_seqs:
- result = benchmark_prefill(
- bs,
- max_seq_len,
- dtype=torch.bfloat16,
- kv_cache_dtype="auto",
- )
- all_results.append(result)
+ dtype = torch.bfloat16
+ quant_dtypes = [
+ # (q_quant_dtype, kv_quant_dtype, o_quant_dtype)
+ (None, None, None),
+ (FP8_DTYPE, FP8_DTYPE, FP8_DTYPE),
+ (FP8_DTYPE, FP8_DTYPE, FP4_DTYPE),
+ ]
+
+ for quant_dtype in quant_dtypes:
+ q_quant_dtype, kv_quant_dtype, o_quant_dtype = quant_dtype
+ q_quant_dtype = q_quant_dtype or dtype
+ kv_quant_dtype = kv_quant_dtype or dtype
+ o_quant_dtype = o_quant_dtype or dtype
+
+ print(
+ f"Running benchmark for q_dtype = {q_quant_dtype}, "
+ f"kv_cache_dtype: {kv_quant_dtype}, "
+ f"output_dtype: {o_quant_dtype}"
+ )
+ print(
+ "\tbatch_size\tmax_seq_len\ttrtllm_mean\ttrtllm_std\tbaseline_mean\t"
+ "baseline_std\tspeedup_percent"
+ )
+ for max_seq_len in max_seq_lens:
+ for bs in batch_sizes:
+ result = benchmark_prefill(
+ dtype=dtype,
+ quant_dtypes=quant_dtype,
+ batch_size=bs,
+ max_seq_len=max_seq_len,
+ )
+ all_results.append(result)
# Write all results to CSV
write_results_to_csv(all_results)
diff --git a/benchmarks/kernels/benchmark_w8a8_block_fp8.py b/benchmarks/kernels/benchmark_w8a8_block_fp8.py
index 4fcdbadd65..98bde9d83c 100644
--- a/benchmarks/kernels/benchmark_w8a8_block_fp8.py
+++ b/benchmarks/kernels/benchmark_w8a8_block_fp8.py
@@ -11,8 +11,8 @@ from datetime import datetime
from typing import Any
import torch
-import tqdm
import triton
+from tqdm import tqdm
from vllm.model_executor.layers.quantization.utils.fp8_utils import (
_w8a8_block_fp8_matmul,
@@ -141,6 +141,7 @@ def get_weight_shapes(tp_size):
# cannot TP
total = [
(512 + 64, 7168),
+ (2112, 7168),
((128 + 64) * 128, 7168),
(128 * (128 + 128), 512),
(7168, 16384),
diff --git a/benchmarks/kernels/weight_shapes.py b/benchmarks/kernels/weight_shapes.py
index a27f02394a..9a057990bd 100644
--- a/benchmarks/kernels/weight_shapes.py
+++ b/benchmarks/kernels/weight_shapes.py
@@ -95,4 +95,10 @@ WEIGHT_SHAPES = {
([2048, 2816], 1),
([1408, 2048], 0),
],
+ "CohereLabs/c4ai-command-a-03-2025": [
+ ([12288, 14336], 1),
+ ([12288, 12288], 0),
+ ([12288, 73728], 1),
+ ([36864, 12288], 0),
+ ],
}
diff --git a/benchmarks/kv_cache/benchmark_block_pool.py b/benchmarks/kv_cache/benchmark_block_pool.py
deleted file mode 100644
index 134551bb61..0000000000
--- a/benchmarks/kv_cache/benchmark_block_pool.py
+++ /dev/null
@@ -1,108 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-import gc
-import time
-from typing import Optional
-
-from tabulate import tabulate
-
-from vllm.utils import FlexibleArgumentParser
-from vllm.v1.core.block_pool import BlockPool
-
-
-class Metric:
- def __init__(self) -> None:
- self.cnt: int = 0
- self.sum_v: int = 0
- self.max_v: Optional[int] = None
-
- def update(self, v: int) -> None:
- self.cnt += 1
- self.sum_v += v
- if self.max_v is None:
- self.max_v = v
- else:
- self.max_v = max(self.max_v, v)
-
- def avg_v(self) -> float:
- return self.sum_v * 1.0 / self.cnt
-
-
-def main(args):
- rows = []
- for allocate_block in args.allocate_blocks:
- # Enforce a GC collect ahead to minimize the impact among runs
- gc.collect()
- block_pool = BlockPool(num_gpu_blocks=args.num_gpu_blocks, enable_caching=True)
-
- get_blocks_metric: Metric = Metric()
- free_blocks_metric: Metric = Metric()
- for _ in range(args.num_iteration):
- t1 = time.monotonic_ns()
- blocks = block_pool.get_new_blocks(allocate_block)
- t2 = time.monotonic_ns()
- block_pool.free_blocks(blocks)
- t3 = time.monotonic_ns()
- get_blocks_metric.update(t2 - t1)
- free_blocks_metric.update(t3 - t2)
-
- if get_blocks_metric.max_v is not None and free_blocks_metric.max_v is not None:
- rows.append(
- [
- get_blocks_metric.cnt,
- args.num_gpu_blocks,
- allocate_block,
- get_blocks_metric.avg_v() / 1000000,
- get_blocks_metric.max_v / 1000000.0,
- free_blocks_metric.avg_v() / 1000000,
- free_blocks_metric.max_v / 1000000.0,
- ]
- )
- else:
- print(
- "No valid metrics found."
- f" {get_blocks_metric.max_v=} {free_blocks_metric.max_v=}"
- )
-
- print(
- tabulate(
- rows,
- headers=[
- "Iterations",
- "Total\nBlocks",
- "Allocated\nBlocks",
- "Get Blocks\nAvg (ms)",
- "Get Blocks\nMax (ms)",
- "Free Blocks\nAvg (ms)",
- "Free Blocks\nMax (ms)",
- ],
- tablefmt="grid",
- floatfmt=".6f",
- )
- )
-
-
-def invoke_main() -> None:
- parser = FlexibleArgumentParser(
- description="Benchmark the performance of BlockPool for KV Cache."
- )
- parser.add_argument("--num-gpu-blocks", type=int, default=100000)
- parser.add_argument(
- "--num-iteration",
- type=int,
- default=1000,
- help="Number of iterations to run to stablize final data readings",
- )
- parser.add_argument(
- "--allocate-blocks",
- type=int,
- nargs="*",
- default=[10, 50, 100, 500, 1000],
- help="Number of blocks to allocate",
- )
- args = parser.parse_args()
- main(args)
-
-
-if __name__ == "__main__":
- invoke_main() # pragma: no cover
diff --git a/benchmarks/multi_turn/README.md b/benchmarks/multi_turn/README.md
new file mode 100644
index 0000000000..7adf97bcf5
--- /dev/null
+++ b/benchmarks/multi_turn/README.md
@@ -0,0 +1,73 @@
+# Benchmark KV Cache Offloading with Multi-Turn Conversations
+
+The requirements (pip) for `benchmark_serving_multi_turn.py` can be found in `requirements.txt`
+
+First start serving your model
+
+```bash
+export MODEL_PATH=/models/meta-llama/Meta-Llama-3.1-8B-Instruct/
+
+vllm serve $MODEL_PATH --served-model-name Llama --disable-log-requests
+```
+
+The variable `MODEL_PATH` should be a path to the model files (e.g. downloaded from huggingface).
+
+## Synthetic Multi-Turn Conversations
+
+Download the following text file (used for generation of synthetic conversations)
+
+```bash
+wget https://www.gutenberg.org/ebooks/1184.txt.utf-8
+mv 1184.txt.utf-8 pg1184.txt
+```
+
+The filename `pg1184.txt` is used in `generate_multi_turn.json` (see `"text_files"`).
+
+But you may use other text files if you prefer (using this specific file is not required).
+
+Then run the benchmarking script
+
+```bash
+export MODEL_PATH=/models/meta-llama/Meta-Llama-3.1-8B-Instruct/
+
+python benchmark_serving_multi_turn.py --model $MODEL_PATH --served-model-name Llama \
+--input-file generate_multi_turn.json --num-clients 2 --max-active-conversations 6
+```
+
+You can edit the file `generate_multi_turn.json` to change the conversation parameters (number of turns, etc.).
+
+If successful, you will see the following output
+
+```bash
+----------------------------------------------------------------------------------------------------
+Statistics summary:
+runtime_sec = 215.810
+requests_per_sec = 0.769
+----------------------------------------------------------------------------------------------------
+ count mean std min 25% 50% 75% 90% 99% max
+ttft_ms 166.0 78.22 67.63 45.91 59.94 62.26 64.43 69.66 353.18 567.54
+tpot_ms 166.0 25.37 0.57 24.40 25.07 25.31 25.50 25.84 27.50 28.05
+latency_ms 166.0 2591.07 326.90 1998.53 2341.62 2573.01 2860.10 3003.50 3268.46 3862.94
+input_num_turns 166.0 7.43 4.57 1.00 3.00 7.00 11.00 13.00 17.00 17.00
+input_num_tokens 166.0 2006.20 893.56 522.00 1247.75 2019.00 2718.00 3233.00 3736.45 3899.00
+output_num_tokens 166.0 100.01 11.80 80.00 91.00 99.00 109.75 116.00 120.00 120.00
+output_num_chunks 166.0 99.01 11.80 79.00 90.00 98.00 108.75 115.00 119.00 119.00
+----------------------------------------------------------------------------------------------------
+```
+
+## ShareGPT Conversations
+
+To run with the ShareGPT data, download the following ShareGPT dataset:
+`https://huggingface.co/datasets/philschmid/sharegpt-raw/blob/main/sharegpt_20230401_clean_lang_split.json`
+
+Use the `convert_sharegpt_to_openai.py` script to convert the dataset to a format supported by `benchmark_serving_multi_turn.py`
+
+```bash
+python convert_sharegpt_to_openai.py sharegpt_20230401_clean_lang_split.json sharegpt_conv_128.json --seed=99 --max-items=128
+```
+
+The script will convert the ShareGPT dataset to a dataset with the standard user/assistant roles.
+
+The flag `--max-items=128` is used to sample 128 conversations from the original dataset (change as needed).
+
+Use the output JSON file `sharegpt_conv_128.json` as the `--input-file` for `benchmark_serving_multi_turn.py`.
diff --git a/benchmarks/multi_turn/bench_dataset.py b/benchmarks/multi_turn/bench_dataset.py
new file mode 100644
index 0000000000..411b89dd23
--- /dev/null
+++ b/benchmarks/multi_turn/bench_dataset.py
@@ -0,0 +1,493 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+from abc import ABC, abstractmethod
+from statistics import mean
+from typing import Any, NamedTuple, Optional, Union
+
+import numpy as np # type: ignore
+import pandas as pd # type: ignore
+from bench_utils import (
+ TEXT_SEPARATOR,
+ Color,
+ logger,
+)
+from transformers import AutoTokenizer # type: ignore
+
+# Conversation ID is a string (e.g: "UzTK34D")
+ConvId = str
+
+# A list of dicts (dicts with keys "id" and "messages")
+ShareGptConversations = list[dict[str, Any]]
+
+# A list of dicts (dicts with keys "role" and "content")
+MessagesList = list[dict[str, str]]
+
+# Map conversation ID to conversation messages
+ConversationsMap = list[ConvId, MessagesList]
+
+
+class Distribution(ABC):
+ @abstractmethod
+ def sample(self, size: int = 1) -> np.ndarray:
+ pass
+
+
+class UniformDistribution(Distribution):
+ def __init__(
+ self,
+ min_val: Union[int, float],
+ max_val: Union[int, float],
+ is_integer: bool = True,
+ ) -> None:
+ self.min_val = min_val
+ self.max_val = max_val
+ self.is_integer = is_integer
+
+ def sample(self, size: int = 1) -> np.ndarray:
+ if self.is_integer:
+ return np.random.randint(
+ int(self.min_val), int(self.max_val + 1), size=size
+ )
+ else:
+ return np.random.uniform(self.min_val, self.max_val, size=size)
+
+ def __repr__(self) -> str:
+ return f"UniformDistribution[{self.min_val}, {self.max_val}]"
+
+
+class ConstantDistribution(Distribution):
+ def __init__(self, value: Union[int, float]) -> None:
+ self.value = value
+ self.max_val = value
+
+ def sample(self, size: int = 1) -> np.ndarray:
+ return np.full(shape=size, fill_value=self.value)
+
+ def __repr__(self) -> str:
+ return f"Constant[{self.value}]"
+
+
+class ZipfDistribution(Distribution):
+ def __init__(self, alpha: float, max_val: Optional[int] = None) -> None:
+ self.alpha = alpha
+ self.max_val = max_val
+
+ def sample(self, size: int = 1) -> np.ndarray:
+ samples = np.random.zipf(self.alpha, size=size)
+ if self.max_val:
+ samples = np.minimum(samples, self.max_val)
+ return samples
+
+ def __repr__(self) -> str:
+ return f"ZipfDistribution[{self.alpha}]"
+
+
+class PoissonDistribution(Distribution):
+ def __init__(self, alpha: float, max_val: Optional[int] = None) -> None:
+ self.alpha = alpha
+ self.max_val = max_val
+
+ def sample(self, size: int = 1) -> np.ndarray:
+ samples = np.random.poisson(self.alpha, size=size)
+ if self.max_val:
+ samples = np.minimum(samples, self.max_val)
+ return samples
+
+ def __repr__(self) -> str:
+ return f"PoissonDistribution[{self.alpha}]"
+
+
+class LognormalDistribution(Distribution):
+ def __init__(
+ self, mean: float, sigma: float, max_val: Optional[int] = None
+ ) -> None:
+ self.mean = mean
+ self.sigma = sigma
+ self.max_val = max_val
+
+ def sample(self, size: int = 1) -> np.ndarray:
+ samples = np.random.lognormal(mean=self.mean, sigma=self.sigma, size=size)
+ if self.max_val:
+ samples = np.minimum(samples, self.max_val)
+
+ return np.round(samples).astype(int)
+
+ def __repr__(self) -> str:
+ return f"LognormalDistribution[{self.mean}, {self.sigma}]"
+
+
+class GenConvArgs(NamedTuple):
+ num_conversations: int
+ text_files: list[str]
+ input_num_turns: Distribution
+ input_common_prefix_num_tokens: Distribution
+ input_prefix_num_tokens: Distribution
+ input_num_tokens: Distribution
+ output_num_tokens: Distribution
+ print_stats: bool
+
+
+def verify_field_exists(
+ conf: dict, field_name: str, section: str, subsection: str
+) -> None:
+ if field_name not in conf:
+ raise ValueError(
+ f"Missing field '{field_name}' in {section=} and {subsection=}"
+ )
+
+
+def get_random_distribution(
+ conf: dict, section: str, subsection: str, optional: bool = False
+) -> Distribution:
+ # section can be "prompt_input" or "prompt_output" (both required)
+ conf = conf[section]
+
+ if optional and subsection not in conf:
+ # Optional subsection, if not found assume the value is always 0
+ return ConstantDistribution(0)
+
+ # subsection can be "num_turns", "num_tokens" or "prefix_num_tokens"
+ if subsection not in conf:
+ raise ValueError(f"Missing subsection {subsection} in section {section}")
+
+ conf = conf[subsection]
+
+ distribution = conf.get("distribution")
+ if distribution is None:
+ raise ValueError(
+ f"Missing field 'distribution' in {section=} and {subsection=}"
+ )
+
+ if distribution == "constant":
+ verify_field_exists(conf, "value", section, subsection)
+ return ConstantDistribution(conf["value"])
+
+ elif distribution == "zipf":
+ verify_field_exists(conf, "alpha", section, subsection)
+ max_val = conf.get("max", None)
+ return ZipfDistribution(conf["alpha"], max_val=max_val)
+
+ elif distribution == "poisson":
+ verify_field_exists(conf, "alpha", section, subsection)
+ max_val = conf.get("max", None)
+ return PoissonDistribution(conf["alpha"], max_val=max_val)
+
+ elif distribution == "lognormal":
+ verify_field_exists(conf, "mean", section, subsection)
+ verify_field_exists(conf, "sigma", section, subsection)
+ max_val = conf.get("max", None)
+ return LognormalDistribution(conf["mean"], conf["sigma"], max_val=max_val)
+
+ elif distribution == "uniform":
+ verify_field_exists(conf, "min", section, subsection)
+ verify_field_exists(conf, "max", section, subsection)
+
+ min_value = conf["min"]
+ max_value = conf["max"]
+
+ assert min_value > 0
+ assert min_value <= max_value
+
+ is_integer = isinstance(min_value, int) and isinstance(max_value, int)
+ return UniformDistribution(min_value, max_value, is_integer)
+ else:
+ raise ValueError(f"Unknown distribution: {distribution}")
+
+
+def parse_input_json_file(conf: dict) -> GenConvArgs:
+ # Validate the input file
+ assert isinstance(conf, dict)
+ required_fields = [
+ "filetype",
+ "num_conversations",
+ "text_files",
+ "prompt_input",
+ "prompt_output",
+ ]
+ for field in required_fields:
+ assert field in conf, f"Missing field {field} in input {conf}"
+
+ assert conf["filetype"] == "generate_conversations"
+
+ assert conf["num_conversations"] > 0, "num_conversations should be larger than zero"
+
+ text_files = conf["text_files"]
+
+ assert isinstance(text_files, list), "Field 'text_files' should be a list"
+ assert len(text_files) > 0, (
+ "Field 'text_files' should be a list with at least one file"
+ )
+
+ # Parse the parameters for the prompt input/output workload
+ input_num_turns = get_random_distribution(conf, "prompt_input", "num_turns")
+ input_num_tokens = get_random_distribution(conf, "prompt_input", "num_tokens")
+ input_common_prefix_num_tokens = get_random_distribution(
+ conf, "prompt_input", "common_prefix_num_tokens", optional=True
+ )
+ input_prefix_num_tokens = get_random_distribution(
+ conf, "prompt_input", "prefix_num_tokens"
+ )
+ output_num_tokens = get_random_distribution(conf, "prompt_output", "num_tokens")
+
+ print_stats: bool = conf.get("print_stats", False)
+ assert isinstance(print_stats, bool), (
+ "Field 'print_stats' should be either 'true' or 'false'"
+ )
+
+ args = GenConvArgs(
+ num_conversations=conf["num_conversations"],
+ text_files=text_files,
+ input_num_turns=input_num_turns,
+ input_common_prefix_num_tokens=input_common_prefix_num_tokens,
+ input_prefix_num_tokens=input_prefix_num_tokens,
+ input_num_tokens=input_num_tokens,
+ output_num_tokens=output_num_tokens,
+ print_stats=print_stats,
+ )
+ return args
+
+
+def print_conv_stats(conversations: ConversationsMap, tokenizer: AutoTokenizer) -> None:
+ # Collect statistics
+ conv_stats: list[dict[Any, Any]] = []
+ req_stats: list[int] = []
+
+ print("\nCollecting statistics...")
+ for messages in conversations.values():
+ # messages is a list of dicts
+ user_tokens: list[int] = []
+ assistant_tokens: list[int] = []
+ request_tokens: list[int] = []
+
+ req_tokens = 0
+ for m in messages:
+ content = m["content"]
+ num_tokens = len(tokenizer(content).input_ids)
+
+ if m["role"] == "user":
+ user_tokens.append(num_tokens)
+ # New user prompt including all chat history
+ req_tokens += num_tokens
+ request_tokens.append(req_tokens)
+
+ elif m["role"] == "assistant":
+ assistant_tokens.append(num_tokens)
+ # Update assistant answer
+ # (will be part of chat history for the next user prompt)
+ req_tokens += num_tokens
+
+ item_stats = {
+ "conversation_turns": len(messages),
+ "user_tokens": mean(user_tokens),
+ "assistant_tokens": mean(assistant_tokens),
+ }
+
+ conv_stats.append(item_stats)
+ req_stats.extend(request_tokens)
+
+ # Print statistics
+ percentiles = [0.25, 0.5, 0.75, 0.9, 0.99]
+
+ print(TEXT_SEPARATOR)
+ print(f"{Color.YELLOW}Conversations statistics:{Color.RESET}")
+ print(TEXT_SEPARATOR)
+ df = pd.DataFrame(conv_stats)
+ print(df.describe(percentiles=percentiles).transpose())
+ print(TEXT_SEPARATOR)
+ print(f"{Color.YELLOW}Request statistics:{Color.RESET}")
+ print(TEXT_SEPARATOR)
+ df = pd.DataFrame(req_stats, columns=["request_tokens"])
+ print(df.describe(percentiles=percentiles).transpose())
+ print(TEXT_SEPARATOR)
+
+
+def generate_conversations(
+ args: GenConvArgs, tokenizer: AutoTokenizer
+) -> ConversationsMap:
+ # Text for all user prompts
+ # (text from the input text files will be appended to this line)
+ base_prompt_text = "Please rewrite the following text and add more content: "
+ base_prompt_token_count = len(
+ tokenizer.encode(base_prompt_text, add_special_tokens=False)
+ )
+
+ logger.info(f"{Color.PURPLE}Generating conversations...{Color.RESET}")
+ logger.info(args)
+
+ list_of_tokens = []
+
+ for filename in args.text_files:
+ # Load text file that will be used to generate prompts
+ with open(filename) as file:
+ data = file.read()
+ tokens_in_file = tokenizer.encode(data, add_special_tokens=False)
+ list_of_tokens.extend(tokens_in_file)
+
+ conversations: ConversationsMap = {}
+ conv_id = 0
+
+ # Generate number of turns for every conversation
+ turn_count: np.ndarray = args.input_num_turns.sample(args.num_conversations)
+
+ # Turn count should be at least 2 (one user prompt and one assistant answer)
+ turn_count = np.maximum(turn_count, 2)
+
+ # Round up to an even number (every user prompt should have an answer)
+ turn_count = turn_count + (turn_count % 2)
+
+ # Generate number of prefix tokens for every conversation
+ conv_prefix_tokens: np.ndarray = args.input_prefix_num_tokens.sample(
+ args.num_conversations
+ )
+
+ # Used to reduce shared text between conversations
+ # (jump/skip over text sections between conversations)
+ base_offset = 0
+
+ # Common prefix size for all conversations (only 1 sample required)
+ common_prefix_text = ""
+ common_prefix_tokens: int = args.input_common_prefix_num_tokens.sample(1)[0]
+ if common_prefix_tokens > 0:
+ # Using "." at the end to separate sentences
+ common_prefix_text = (
+ tokenizer.decode(list_of_tokens[: common_prefix_tokens - 2]) + "."
+ )
+ base_offset += common_prefix_tokens
+
+ for conv_id in range(args.num_conversations):
+ # Generate a single conversation
+ messages: MessagesList = []
+
+ nturns = turn_count[conv_id]
+
+ # User prompt token count per turn (with lower limit)
+ input_token_count: np.ndarray = args.input_num_tokens.sample(nturns)
+ input_token_count = np.maximum(input_token_count, base_prompt_token_count)
+
+ # Assistant answer token count per turn (with lower limit)
+ output_token_count: np.ndarray = args.output_num_tokens.sample(nturns)
+ output_token_count = np.maximum(output_token_count, 1)
+
+ user_turn = True
+ for turn_id in range(nturns):
+ if user_turn:
+ role = "user"
+ num_tokens = input_token_count[turn_id]
+
+ # Generate the user prompt,
+ # use a unique prefix (the conv_id) for each conversation
+ # (to avoid shared prefix between conversations)
+ content = f"{conv_id} is a nice number... "
+
+ if len(common_prefix_text) > 0 and turn_id == 0:
+ content = common_prefix_text + content
+
+ # Update the number of tokens left for the content
+ num_tokens -= len(tokenizer.encode(content, add_special_tokens=False))
+
+ if turn_id == 0:
+ prefix_num_tokens = conv_prefix_tokens[conv_id]
+ if prefix_num_tokens > 0:
+ # Add prefix text (context) to the first turn
+ start_offset = base_offset
+ end_offset = start_offset + prefix_num_tokens
+ assert len(list_of_tokens) > end_offset, (
+ "Not enough input text to generate "
+ f"{prefix_num_tokens} tokens for the "
+ f"prefix text ({start_offset=}, {end_offset=})"
+ )
+
+ content += f"{conv_id}, " + tokenizer.decode(
+ list_of_tokens[start_offset:end_offset]
+ )
+ base_offset += prefix_num_tokens
+
+ # Add the actual user prompt/question after the prefix text
+ content += base_prompt_text
+ num_tokens -= base_prompt_token_count
+
+ if num_tokens > 0:
+ # Add text from the input file (to reach the desired token count)
+ start_offset = base_offset + turn_id * input_token_count.max()
+ end_offset = start_offset + num_tokens
+ assert len(list_of_tokens) > end_offset, (
+ f"Not enough input text to generate {num_tokens} tokens "
+ f"for the prompt ({start_offset=}, {end_offset=})"
+ )
+
+ # Convert tokens back to text
+ content += tokenizer.decode(list_of_tokens[start_offset:end_offset])
+ else:
+ role = "assistant"
+ # This content will not be used as input to the LLM server
+ # (actual answers will be used instead).
+ # Content is only required to determine the min_tokens/max_tokens
+ # (inputs to the LLM server).
+ num_tokens = output_token_count[turn_id]
+ assert len(list_of_tokens) > num_tokens, (
+ f"Not enough input text to generate {num_tokens} "
+ "tokens for assistant content"
+ )
+ content = tokenizer.decode(list_of_tokens[:num_tokens])
+
+ # Append the user/assistant message to the list of messages
+ messages.append({"role": role, "content": content})
+ user_turn = not user_turn
+
+ # Add the new conversation
+ conversations[f"CONV_ID_{conv_id}"] = messages
+
+ # Increase base offset for the next conversation
+ base_offset += nturns
+
+ if args.print_stats:
+ print_conv_stats(conversations, tokenizer)
+
+ return conversations
+
+
+def conversations_list_to_dict(input_list: ShareGptConversations) -> ConversationsMap:
+ conversations: ConversationsMap = {}
+
+ for item in input_list:
+ conv_id: str = item["id"]
+ assert isinstance(conv_id, str)
+
+ assert conv_id not in conversations, (
+ f"Conversation ID {conv_id} found more than once in the input"
+ )
+
+ messages: MessagesList = item["messages"]
+ assert isinstance(messages, list), (
+ f"Conversation messages should be a list (ID: {conv_id})"
+ )
+ assert len(messages) > 0, f"Conversation with no messages (ID: {conv_id})"
+
+ conversations[conv_id] = messages
+
+ logger.info(f"Using {len(conversations)} unique conversations (IDs)")
+ assert len(conversations) == len(input_list)
+
+ # Print statistics about the selected conversations
+ stats: list[dict[str, Any]] = []
+ for conv_data in conversations.values():
+ stats.append({"num_turns": len(conv_data)})
+
+ print(TEXT_SEPARATOR)
+ print(f"{Color.YELLOW}Conversations statistics:{Color.RESET}")
+ print(TEXT_SEPARATOR)
+ percentiles = [0.25, 0.5, 0.75, 0.9, 0.99, 0.999, 0.9999]
+ conv_stats = pd.DataFrame(stats).describe(percentiles=percentiles)
+ print(conv_stats.transpose())
+ print(TEXT_SEPARATOR)
+
+ return conversations
+
+
+def conversations_dict_to_list(input_dict: ConversationsMap) -> ShareGptConversations:
+ output: ShareGptConversations = []
+ for conv_id, conv_data in input_dict.items():
+ new_item = {"id": conv_id, "messages": conv_data}
+ output.append(new_item)
+
+ return output
diff --git a/benchmarks/multi_turn/bench_utils.py b/benchmarks/multi_turn/bench_utils.py
new file mode 100644
index 0000000000..e959a4be71
--- /dev/null
+++ b/benchmarks/multi_turn/bench_utils.py
@@ -0,0 +1,28 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import logging
+from enum import Enum
+
+
+class Color(Enum):
+ RED = "\033[91m"
+ GREEN = "\033[92m"
+ BLUE = "\033[94m"
+ PURPLE = "\033[95m"
+ CYAN = "\033[96m"
+ YELLOW = "\033[93m"
+ RESET = "\033[0m"
+
+ def __str__(self):
+ return self.value
+
+
+TEXT_SEPARATOR = "-" * 100
+
+# Configure the logger
+logging.basicConfig(
+ level=logging.INFO,
+ format="%(asctime)s [%(levelname)s] - %(message)s",
+ datefmt="%d-%m-%Y %H:%M:%S",
+)
+logger = logging.getLogger(__name__)
diff --git a/benchmarks/multi_turn/benchmark_serving_multi_turn.py b/benchmarks/multi_turn/benchmark_serving_multi_turn.py
new file mode 100644
index 0000000000..66d85eaf51
--- /dev/null
+++ b/benchmarks/multi_turn/benchmark_serving_multi_turn.py
@@ -0,0 +1,1569 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import argparse
+import asyncio
+import json
+import logging
+import multiprocessing as mp
+import os
+import random
+import time
+from collections import Counter, deque
+from datetime import datetime
+from enum import Enum
+from http import HTTPStatus
+from statistics import mean
+from typing import NamedTuple, Optional, Union
+
+import aiohttp # type: ignore
+import numpy as np # type: ignore
+import pandas as pd # type: ignore
+from bench_dataset import (
+ ConversationsMap,
+ ConvId,
+ GenConvArgs,
+ MessagesList,
+ ShareGptConversations,
+ conversations_dict_to_list,
+ conversations_list_to_dict,
+ generate_conversations,
+ parse_input_json_file,
+)
+from bench_utils import TEXT_SEPARATOR, Color, logger
+from transformers import AutoTokenizer # type: ignore
+
+NUM_TOKENS_FROM_DATASET = 0
+TERM_SIGNAL = None
+
+
+class ConversationSampling(str, Enum):
+ ROUND_ROBIN = "round_robin"
+ RANDOM = "random"
+
+ def __str__(self):
+ return self.value
+
+
+class ClientArgs(NamedTuple):
+ seed: int
+ max_num_requests: Optional[int]
+ skip_first_turn: bool
+ max_turns: Optional[int]
+ max_active_conversations: int
+ verbose: bool
+ print_content: bool
+ verify_output: bool
+ conversation_sampling: ConversationSampling
+ request_rate: float
+
+
+class RequestArgs(NamedTuple):
+ chat_url: str
+ model: str
+ stream: bool
+ limit_min_tokens: int # Use negative value for no limit
+ limit_max_tokens: int # Use negative value for no limit
+
+
+class BenchmarkArgs(NamedTuple):
+ url: str
+ num_clients: int
+ early_stop: bool
+
+
+class ServerResponse(NamedTuple):
+ valid: bool
+ ttft_ms: float # time to first chunk
+ tpot_ms: float # time per output chunk (one or more tokens)
+ latency_ms: float
+ start_time_ms: float
+ first_chunk: str # first chunk of the content
+ content: str # includes the first_chunk
+ num_chunks: int
+
+ def __str__(self) -> str:
+ return f"ttft_ms {self.ttft_ms:.2f}, tpot_ms {self.tpot_ms:.2f}, latency_ms {self.latency_ms:.2f}" # noqa: E501
+
+
+class RequestStats(NamedTuple):
+ ttft_ms: float
+ tpot_ms: float
+ latency_ms: float
+ start_time_ms: float
+ input_num_turns: int
+ input_num_tokens: int
+ output_num_tokens: int
+ output_num_chunks: int
+ output_num_first_chunk_tokens: int
+ approx_cached_percent: float
+ conversation_id: str
+ client_id: int
+
+ def __str__(self) -> str:
+ return (
+ f"ttft_ms {self.ttft_ms:.2f}, tpot_ms {self.tpot_ms:.2f}, latency_ms {self.latency_ms:.2f}, input_num_tokens {self.input_num_tokens}, " # noqa: E501
+ f"output_num_tokens {self.output_num_tokens} ({self.output_num_chunks} chunks, {self.output_num_first_chunk_tokens} tokens in first chunk), " # noqa: E501
+ f"approx_cached_percent {self.approx_cached_percent:.2f}%"
+ )
+
+
+class MetricStats:
+ def __init__(self) -> None:
+ self.min: Optional[float] = None
+ self.max: Optional[float] = None
+ self.avg: Optional[float] = None
+ self.sum = 0.0
+ self.count = 0
+
+ def update(self, value: float) -> None:
+ if self.min is None:
+ self.min = value
+ else:
+ self.min = min(self.min, value)
+
+ if self.max is None:
+ self.max = value
+ else:
+ self.max = max(self.max, value)
+
+ self.sum += value
+ self.count += 1
+ self.avg = self.sum / self.count
+
+ def __repr__(self) -> str:
+ if self.count == 0:
+ return "no data"
+ return f"avg: {self.avg:>10.3f}, min: {self.min:>10.3f}, max: {self.max:>10.3f}"
+
+
+class MovingAverage:
+ def __init__(self, window_size: int) -> None:
+ self.window_size = window_size
+ self.window = np.zeros(window_size)
+ self.index = 0
+ self.sum = 0.0
+ self.count = 0
+ self.avg: Optional[float] = None
+
+ def update(self, new_value: float) -> None:
+ if self.count < self.window_size:
+ # Filling up the window
+ self.sum += new_value
+ self.window[self.count] = new_value
+ self.count += 1
+ else:
+ # Window is full, start replacing old values
+ old_value = self.window[self.index]
+ self.sum = self.sum - old_value + new_value
+ self.window[self.index] = new_value
+ self.index = (self.index + 1) % self.window_size
+
+ self.avg = self.sum / self.count
+
+ def __repr__(self) -> str:
+ if self.count == 0:
+ return "no data"
+ return f"avg: {self.avg:>10.3f} ({self.count} samples)"
+
+
+class DebugStats:
+ def __init__(self, logger: logging.Logger, window_size: int) -> None:
+ self.logger = logger
+ self.metrics: dict[str, Union[MovingAverage, MetricStats]] = {
+ "moving_avg_ttft_ms": MovingAverage(window_size),
+ "moving_avg_tpot_ms": MovingAverage(window_size),
+ "ttft_ms": MetricStats(),
+ "tpot_ms": MetricStats(),
+ "latency_ms": MetricStats(),
+ "input_num_turns": MetricStats(),
+ "input_num_tokens": MetricStats(),
+ "output_num_tokens": MetricStats(),
+ }
+
+ def update(self, data: RequestStats) -> None:
+ self.metrics["ttft_ms"].update(data.ttft_ms)
+ self.metrics["moving_avg_ttft_ms"].update(data.ttft_ms)
+ self.metrics["tpot_ms"].update(data.tpot_ms)
+ self.metrics["moving_avg_tpot_ms"].update(data.tpot_ms)
+ self.metrics["latency_ms"].update(data.latency_ms)
+ self.metrics["input_num_turns"].update(data.input_num_turns)
+ self.metrics["input_num_tokens"].update(data.input_num_tokens)
+ self.metrics["output_num_tokens"].update(data.output_num_tokens)
+
+ def print(self) -> None:
+ self.logger.info("-" * 50)
+ for k, v in self.metrics.items():
+ kv_info = f"[{k:25}] {v}"
+ self.logger.info(kv_info)
+ self.logger.info("-" * 50)
+
+
+# Must support Python 3.8, we can't use str.removeprefix(prefix)
+# introduced in Python 3.9
+def remove_prefix(text: str, prefix: str) -> str:
+ if text.startswith(prefix):
+ return text[len(prefix) :]
+ return text
+
+
+def nanosec_to_millisec(value: float) -> float:
+ return value / 1000000.0
+
+
+def nanosec_to_sec(value: float) -> float:
+ return value / 1000000000.0
+
+
+async def send_request(
+ session: aiohttp.ClientSession,
+ messages: list[dict[str, str]],
+ chat_url: str,
+ model: str,
+ stream: bool = True,
+ min_tokens: Optional[int] = None,
+ max_tokens: Optional[int] = None,
+) -> ServerResponse:
+ payload = {
+ "model": model,
+ "messages": messages,
+ "seed": 0,
+ "temperature": 0.0,
+ }
+
+ if stream:
+ payload["stream"] = True
+ payload["stream_options"] = {"include_usage": False}
+
+ if min_tokens is not None:
+ payload["min_tokens"] = min_tokens
+
+ if max_tokens is not None:
+ payload["max_tokens"] = max_tokens
+
+ headers = {"Content-Type": "application/json"}
+
+ # Calculate the timeout for the request
+ timeout_sec = 120
+ if max_tokens is not None:
+ # Assume TPOT of 200ms and use max_tokens to determine timeout
+ timeout_sec = max(timeout_sec, int(max_tokens * 0.2))
+ timeout = aiohttp.ClientTimeout(total=timeout_sec)
+
+ valid_response = True
+ ttft: Optional[float] = None
+ chunk_delay: list[int] = []
+ latency: Optional[float] = None
+ first_chunk = ""
+ generated_text = ""
+
+ start_time: int = time.perf_counter_ns()
+ most_recent_timestamp: int = start_time
+
+ async with session.post(
+ url=chat_url, json=payload, headers=headers, timeout=timeout
+ ) as response:
+ http_status = HTTPStatus(response.status)
+ if http_status == HTTPStatus.OK:
+ async for chunk_bytes in response.content:
+ chunk_bytes = chunk_bytes.strip()
+ if not chunk_bytes:
+ continue
+
+ chunk = remove_prefix(chunk_bytes.decode("utf-8"), "data: ")
+ if chunk == "[DONE]":
+ # End of stream
+ latency = time.perf_counter_ns() - start_time
+ elif stream is False:
+ data = json.loads(chunk)
+ message = data["choices"][0]["message"]
+ assert message["role"] == "assistant"
+ generated_text += message["content"]
+ else:
+ timestamp: int = time.perf_counter_ns()
+ data = json.loads(chunk)
+
+ # Delta is the new content/text/data
+ delta = data["choices"][0]["delta"]
+ if delta.get("content", None):
+ if ttft is None:
+ # First token
+ first_token_time = time.perf_counter_ns()
+ ttft = first_token_time - start_time
+ first_chunk = delta["content"]
+ else:
+ # Decoding phase
+ chunk_delay.append(timestamp - most_recent_timestamp)
+
+ generated_text += delta["content"]
+
+ most_recent_timestamp = timestamp
+ else:
+ valid_response = False
+ content = await response.text()
+ logger.warning(
+ f"{Color.YELLOW}Received HTTP status {http_status.value} "
+ f"({http_status.phrase}): {content}{Color.RESET}"
+ )
+
+ if latency is None:
+ latency = -1.0
+ if valid_response:
+ # Streaming is disabled, latency was not set
+ latency = time.perf_counter_ns() - start_time
+
+ if ttft is None:
+ # The response was a single chunk
+ ttft = latency
+
+ # Each chunk may include more than one token
+ tpot: float = mean(chunk_delay) if len(chunk_delay) > 0 else 0.0
+ num_chunks: int = len(chunk_delay)
+
+ sr = ServerResponse(
+ valid=valid_response,
+ ttft_ms=nanosec_to_millisec(ttft) if ttft > 0.0 else -1.0,
+ tpot_ms=nanosec_to_millisec(tpot),
+ latency_ms=nanosec_to_millisec(latency),
+ start_time_ms=nanosec_to_millisec(start_time),
+ first_chunk=first_chunk,
+ content=generated_text,
+ num_chunks=num_chunks,
+ )
+ return sr
+
+
+def get_short_string(input: str) -> str:
+ n = 20
+ if len(input) < 400:
+ return input
+
+ return f"{input[:n]}...{input[-n:]}"
+
+
+def get_token_count(tokenizer: AutoTokenizer, text: str) -> int:
+ return len(tokenizer(text, add_special_tokens=False).input_ids)
+
+
+def get_messages_token_count(
+ tokenizer: AutoTokenizer, messages: list[dict[str, str]]
+) -> int:
+ token_count = 0
+ for m in messages:
+ token_count += get_token_count(tokenizer, m["content"])
+
+ return token_count
+
+
+async def send_turn(
+ session: aiohttp.ClientSession,
+ client_id: int,
+ conv_id: str,
+ conversation_messages: MessagesList,
+ messages_to_use: int,
+ tokenizer: AutoTokenizer,
+ req_args: RequestArgs,
+ verbose: bool,
+ verify_output: bool,
+) -> Optional[RequestStats]:
+ assert messages_to_use > 0
+ assert messages_to_use <= len(conversation_messages)
+
+ messages = conversation_messages[:messages_to_use]
+
+ # Index of the next message (the role should be "user")
+ index = messages_to_use - 1
+
+ # Verify that the message has only two keys, "role" and "content"
+ assert len(messages[index].keys()) == 2
+ assert "role" in messages[index] and "content" in messages[index]
+ assert messages[index]["role"] == "user", (
+ f"Failed on conversation ID {conv_id}, message role should be user"
+ )
+
+ if verbose:
+ print(
+ f"{Color.CYAN}Messages (conversation ID {conv_id},"
+ f" {len(messages)} turns):{Color.RESET}",
+ messages,
+ )
+
+ # None means that there is no upper/lower limit for the output token count
+ min_tokens = None if req_args.limit_min_tokens < 0 else req_args.limit_min_tokens
+ max_tokens = None if req_args.limit_max_tokens < 0 else req_args.limit_max_tokens
+
+ if len(conversation_messages) > messages_to_use:
+ # The conversation contains an assistant answer for the next user prompt
+ if (
+ min_tokens == NUM_TOKENS_FROM_DATASET
+ or max_tokens == NUM_TOKENS_FROM_DATASET
+ ):
+ # Compute number of tokens in the answer (from the input conversation)
+ assistant_answer = conversation_messages[messages_to_use]
+ answer_num_tokens = get_token_count(tokenizer, assistant_answer["content"])
+ assert assistant_answer["role"] == "assistant"
+
+ if min_tokens == NUM_TOKENS_FROM_DATASET:
+ min_tokens = max(1, answer_num_tokens)
+
+ if max_tokens == NUM_TOKENS_FROM_DATASET:
+ max_tokens = max(1, answer_num_tokens)
+
+ # Send the current conversation to LLM and get a response
+ response: ServerResponse = await send_request(
+ session,
+ messages,
+ req_args.chat_url,
+ req_args.model,
+ req_args.stream,
+ min_tokens,
+ max_tokens,
+ )
+
+ if response.valid is False:
+ # Request failed
+ return None
+
+ # Compute number of tokens in input / output
+ input_num_tokens = get_messages_token_count(tokenizer, messages)
+
+ # Num tokens in the user's last question
+ question_num_tokens = get_token_count(tokenizer, messages[index]["content"])
+
+ # Num tokens in the history/context of the question
+ assert input_num_tokens >= question_num_tokens
+ history_num_tokens = input_num_tokens - question_num_tokens
+
+ # Num tokens in the LLM's answer (first chunk and full answer)
+ first_chunk_tokens = get_token_count(tokenizer, response.first_chunk)
+
+ output_content = response.content
+ output_num_tokens = get_token_count(tokenizer, output_content)
+
+ # Prefix caching approximated cached percent
+ approx_cached_percent = (
+ 100.0 * (history_num_tokens / input_num_tokens) if input_num_tokens > 0 else 0.0
+ )
+
+ # Compute the correct TTFT and TPOT (based on tokens and not chunks).
+ # Required because multiple output tokens may be bundled in a single chunk.
+ if output_num_tokens > 1 and output_num_tokens > first_chunk_tokens:
+ # More than one token and more than one chunk in the output
+ decode_ms = response.latency_ms - response.ttft_ms
+ decode_num_tokens = output_num_tokens - first_chunk_tokens
+ tpot_ms = decode_ms / decode_num_tokens
+ else:
+ # In this case: output_num_tokens == first_chunk_tokens
+ # Output was a single chunk (output_num_tokens > 1)
+ # or even a single token (output_num_tokens == 1)
+ tpot_ms = 0.0
+
+ if first_chunk_tokens > 1:
+ # First chunk had multiple tokens, adjust TTFT for a single token
+ delta_ms = (first_chunk_tokens - 1) * tpot_ms
+ ttft_ms = max(0.1, response.ttft_ms - delta_ms)
+ else:
+ # First chunk had only one token
+ ttft_ms = response.ttft_ms
+
+ rs = RequestStats(
+ ttft_ms=ttft_ms,
+ tpot_ms=tpot_ms,
+ latency_ms=response.latency_ms,
+ start_time_ms=response.start_time_ms,
+ input_num_turns=len(messages),
+ input_num_tokens=input_num_tokens,
+ output_num_tokens=output_num_tokens,
+ output_num_chunks=response.num_chunks,
+ output_num_first_chunk_tokens=first_chunk_tokens,
+ approx_cached_percent=approx_cached_percent,
+ conversation_id=conv_id,
+ client_id=client_id,
+ )
+
+ if verbose:
+ print(
+ f"\n{Color.YELLOW}Response ({output_num_tokens} tokens):{Color.RESET}",
+ output_content,
+ )
+ print(f"{Color.YELLOW}Response metrics: {rs}{Color.RESET}")
+ print("-" * 70)
+
+ # Save the LLM's answer (will be used as part of the context for the next user turn)
+ answer_index = messages_to_use
+ if len(conversation_messages) > answer_index:
+ assert conversation_messages[answer_index]["role"] == "assistant", (
+ f"Failed on conversation ID {conv_id}, message role should be assistant"
+ )
+
+ orig_content = conversation_messages[answer_index]["content"]
+ if verify_output:
+ # Compare the new answer to the answer from the input file
+ debug_info = (
+ f"LLM/dataset answers do not match ({conv_id}):"
+ f"\n'{get_short_string(output_content)}' (len: {len(output_content)}),"
+ f"\n'{get_short_string(orig_content)}' (len: {len(orig_content)})"
+ )
+ if orig_content != output_content:
+ raise ValueError(debug_info)
+
+ # Update the answer
+ conversation_messages[answer_index]["content"] = output_content
+ else:
+ # A user prompt that has no answer, add the answer as a new message
+ new_answer = {"role": "assistant", "content": output_content}
+ conversation_messages.append(new_answer)
+
+ return rs
+
+
+async def poisson_sleep(request_rate: float, verbose: bool = False) -> None:
+ # Generate a random time interval from the Poisson distribution
+ assert request_rate > 0
+
+ interval = np.random.exponential(1.0 / request_rate)
+ if verbose:
+ logger.info(f"Sleeping for {interval:.3f} seconds...")
+ await asyncio.sleep(interval)
+
+
+async def client_main(
+ args: ClientArgs,
+ req_args: RequestArgs,
+ client_id: int,
+ tokenizer: AutoTokenizer,
+ stop_event: mp.Event, # type: ignore
+ task_queue: mp.Queue,
+ result_queue: mp.Queue,
+ conv_queue: mp.Queue,
+) -> None:
+ logger.info(
+ f"{Color.CYAN}Started client {client_id}: max_num_requests={args.max_num_requests}, max_active_conversations={args.max_active_conversations}{Color.RESET}" # noqa: E501
+ )
+
+ random.seed(args.seed)
+ np.random.seed(args.seed)
+
+ # Active conversations
+ active_convs: ConversationsMap = {}
+ conv_id_queue: deque = deque(maxlen=args.max_active_conversations)
+
+ # Keep track of how many messages have been used for each conversation
+ turns_count: Counter = Counter()
+ num_successes = 0
+ num_failures = 0
+
+ # Track the timestamp (time.perf_counter())
+ # of the last turn per conversation (only for debug)
+ time_of_last_turn: dict[ConvId, float] = {}
+
+ # Flag that indicates that there are no new tasks (conversations) for the client
+ task_queue_empty = False
+
+ async with aiohttp.ClientSession() as session:
+ # Print progress
+
+ while task_queue_empty is False:
+ result = None
+
+ if (
+ args.max_num_requests
+ and num_successes + num_failures == args.max_num_requests
+ ):
+ logger.info(
+ f"{Color.YELLOW}Client {client_id} reached "
+ f"request limit{Color.RESET}"
+ )
+ break
+
+ if stop_event.is_set(): # type: ignore
+ logger.info(
+ f"{Color.YELLOW}Client {client_id} received "
+ f"a termination signal{Color.RESET}"
+ )
+ break
+
+ while (
+ len(active_convs) < args.max_active_conversations
+ and task_queue_empty is False
+ ):
+ # Get a new conversation from the task queue
+ conv_id, messages = task_queue.get()
+
+ if conv_id is TERM_SIGNAL:
+ task_queue_empty = True
+ break
+
+ if args.skip_first_turn:
+ # Skip the first turn (both user and assistant),
+ # relevant if warmup was enabled.
+ # Default turns_count[conv_id] will be zero if conv_id
+ # was never inserted/updated in turns_count.
+ turns_count[conv_id] += 2
+
+ if turns_count[conv_id] < len(messages):
+ # Add new conversation
+ active_convs[conv_id] = messages
+ conv_id_queue.append(conv_id)
+
+ if args.verbose:
+ logger.info(
+ f"{Color.GREEN}Client {client_id} will use conversation ID {conv_id} (active conversations {len(active_convs)}){Color.RESET}" # noqa: E501
+ )
+
+ elif args.verbose:
+ # No more messages (conversation finished during the warmup)
+ logger.info(
+ f"{Color.YELLOW}Client {client_id} will not use conversation ID {conv_id} (all {len(messages)} messages already sent){Color.RESET}" # noqa: E501
+ )
+
+ if len(active_convs) == 0 or task_queue_empty:
+ logger.info(
+ f"{Color.YELLOW}Client {client_id} has no more work{Color.RESET}"
+ )
+ break
+
+ # Pick an active conversation for the next request
+ if args.conversation_sampling == ConversationSampling.ROUND_ROBIN:
+ conv_id = conv_id_queue.pop()
+ else:
+ # ConversationSampling.RANDOM
+ active_ids = list(active_convs.keys())
+ conv_id = random.choice(active_ids)
+
+ messages = active_convs[conv_id]
+ assert isinstance(messages, list) and len(messages) > 0
+
+ # Update the amount of messages to use
+ turns_count[conv_id] += 1
+ current_turn = turns_count[conv_id]
+
+ assert current_turn < len(messages), (
+ f"Turn number {current_turn} is invalid for conversation ID {conv_id}"
+ f" that has only {len(messages)} messages"
+ )
+
+ if args.verbose:
+ curr_time_sec: float = time.perf_counter()
+ time_since_last_turn: Union[str, float] = "N/A"
+ if conv_id in time_of_last_turn:
+ time_since_last_turn = round(
+ curr_time_sec - time_of_last_turn[conv_id], 3
+ )
+ logger.info(
+ f"Client {client_id} using conversation ID {conv_id} (turn: {current_turn}, time since last turn [sec]: {time_since_last_turn})" # noqa: E501
+ )
+ time_of_last_turn[conv_id] = curr_time_sec
+
+ success = True
+ try:
+ result = await send_turn(
+ session,
+ client_id,
+ conv_id,
+ messages,
+ current_turn,
+ tokenizer,
+ req_args,
+ args.print_content,
+ args.verify_output,
+ )
+ if result is not None:
+ result_queue.put(result)
+ else:
+ # None means that the request failed,
+ # and should not be added to the statistics.
+ success = False
+ num_failures += 1
+
+ logger.warning(
+ f"{Color.YELLOW}Client {client_id} - Request rejected during conversation ID {conv_id} (turn: {current_turn}){Color.RESET}" # noqa: E501
+ )
+
+ # Remove the conversation (should not be used again)
+ active_convs.pop(conv_id)
+
+ except asyncio.exceptions.TimeoutError:
+ num_failures += 1
+ logger.exception(
+ f"{Color.RED}Client {client_id} - Timeout during conversation ID {conv_id} (turn: {current_turn}){Color.RESET}" # noqa: E501
+ )
+ break # Exit gracefully instead of raising an error
+
+ except Exception:
+ num_failures += 1
+ logger.exception(
+ f"{Color.RED}Client {client_id} - Exception during conversation ID {conv_id} (turn: {current_turn}){Color.RESET}" # noqa: E501
+ )
+ break # Exit gracefully instead of raising an error
+
+ if success:
+ num_successes += 1
+
+ # Update the turns counter to include the LLM response
+ # The LLM response will be used as context for the next user turn
+ turns_count[conv_id] += 1
+
+ max_turns = len(messages)
+ if args.max_turns is not None:
+ # Limit the number of turns in the conversation
+ max_turns = min(args.max_turns, max_turns)
+
+ if turns_count[conv_id] >= max_turns:
+ # Conversation has no more turns (no longer active)
+ # save the updated conversation (with the LLM server's answer)
+ conv_queue.put((conv_id, active_convs.pop(conv_id)))
+ if args.verbose:
+ logger.info(
+ f"{Color.GREEN}Client {client_id} finished "
+ f"conversation ID {conv_id}{Color.RESET}"
+ )
+ else:
+ # Conversation is not finished, insert it at the back of the queue
+ conv_id_queue.appendleft(conv_id)
+
+ # Sleep between requests (if lambda is positive)
+ if args.request_rate > 0:
+ await poisson_sleep(args.request_rate, args.verbose)
+
+ # Send indication that the client is done
+ conv_queue.put((TERM_SIGNAL, TERM_SIGNAL))
+
+ logger.info(
+ f"{Color.CYAN}Client {client_id} is done "
+ f"({num_successes=}, {num_failures=}){Color.RESET}"
+ )
+
+
+def worker_function(
+ client_id: int,
+ tokenizer: AutoTokenizer,
+ client_args: ClientArgs,
+ req_args: RequestArgs,
+ stop_event: mp.Event, # type: ignore
+ task_queue: mp.Queue,
+ result_queue: mp.Queue,
+ conv_queue: mp.Queue,
+) -> None:
+ asyncio.run(
+ client_main(
+ client_args,
+ req_args,
+ client_id,
+ tokenizer,
+ stop_event,
+ task_queue,
+ result_queue,
+ conv_queue,
+ )
+ )
+
+
+def get_client_config(
+ args: argparse.Namespace, input_conv: ConversationsMap
+) -> tuple[ClientArgs, RequestArgs]:
+ if args.num_clients < 1:
+ raise ValueError("Number of clients must be a positive number")
+
+ if len(input_conv) < args.num_clients:
+ raise ValueError(
+ "Number of conversations must be equal or larger than the number of clients"
+ )
+
+ max_req_per_client: Optional[int] = None
+ if args.max_num_requests is not None:
+ # Max number of requests per client
+ req_per_client = args.max_num_requests // args.num_clients
+ if req_per_client < 1:
+ raise ValueError("Number of requests should be at least one per client")
+ max_req_per_client = req_per_client
+
+ max_active_conversations = args.max_active_conversations
+ if max_active_conversations is None:
+ # Each client will have only one active conversation at a time
+ max_active_conversations = args.num_clients
+
+ if max_active_conversations > len(input_conv):
+ raise ValueError(
+ f"Max active conversations {max_active_conversations} "
+ "must be equal or less than the total number of conversations"
+ )
+
+ # Max number of active conversations per client
+ max_active_conv_per_client = max_active_conversations // args.num_clients
+ if max_active_conv_per_client < 1:
+ raise ValueError(
+ f"Max active conversations {max_active_conversations} "
+ "must be equal or greater than the number of clients"
+ )
+
+ # Skip the first user turn (as part of the warmup)
+ skip_first_turn = args.warmup_step
+
+ # Common arguments for all clients
+ client_args = ClientArgs(
+ seed=args.seed,
+ max_num_requests=max_req_per_client,
+ skip_first_turn=skip_first_turn,
+ max_turns=args.max_turns,
+ max_active_conversations=max_active_conv_per_client,
+ verbose=args.verbose,
+ print_content=args.print_content,
+ verify_output=args.verify_output,
+ conversation_sampling=args.conversation_sampling,
+ request_rate=args.request_rate,
+ )
+
+ if args.limit_min_tokens > 0 or args.limit_max_tokens > 0:
+ if args.limit_min_tokens < 1 or args.limit_max_tokens < 1:
+ raise ValueError(
+ "Invalid min/max tokens limits (both limits should be provided)"
+ )
+ if args.limit_min_tokens > args.limit_max_tokens:
+ raise ValueError(
+ "Invalid min/max tokens limits (min should not be larger than max)"
+ )
+
+ # Arguments for API requests
+ chat_url = f"{args.url}/v1/chat/completions"
+ model_name = args.served_model_name if args.served_model_name else args.model
+
+ req_args = RequestArgs(
+ chat_url=chat_url,
+ model=model_name,
+ stream=not args.no_stream,
+ limit_min_tokens=args.limit_min_tokens,
+ limit_max_tokens=args.limit_max_tokens,
+ )
+
+ return client_args, req_args
+
+
+async def main_mp(
+ client_args: ClientArgs,
+ req_args: RequestArgs,
+ bench_args: BenchmarkArgs,
+ tokenizer: AutoTokenizer,
+ input_conv: ConversationsMap,
+) -> tuple[ConversationsMap, list[RequestStats]]:
+ # An event that will trigger graceful termination of all the clients
+ stop_event = mp.Event()
+
+ # Queue for input conversations (from the input file/dataset)
+ task_queue: mp.Queue = mp.Queue()
+
+ # Queue for client measurements (TTFT, TPOT, etc. for each request)
+ result_queue: mp.Queue = mp.Queue()
+
+ # Queue for output conversations (with the LLM answers, sent by the server)
+ conv_queue: mp.Queue = mp.Queue()
+ output_conv: ConversationsMap = {}
+ client_metrics: list[RequestStats] = []
+
+ # Start all clients
+ start_time = time.perf_counter_ns()
+ logger.info(f"{Color.GREEN}Starting {bench_args.num_clients} clients{Color.RESET}")
+
+ clients = []
+ for client_id in range(bench_args.num_clients):
+ client = mp.Process(
+ name=f"client_{client_id}",
+ target=worker_function,
+ args=(
+ client_id,
+ tokenizer,
+ client_args,
+ req_args,
+ stop_event,
+ task_queue,
+ result_queue,
+ conv_queue,
+ ),
+ )
+ clients.append(client)
+ client.start()
+
+ # Submit all the input conversations as tasks for the clients
+ for conv_id, messages in input_conv.items():
+ task_queue.put((conv_id, messages))
+
+ # Add termination signals for clients
+ for _ in range(bench_args.num_clients):
+ task_queue.put((TERM_SIGNAL, TERM_SIGNAL))
+
+ # Collect the updated conversations from all clients
+ num_clients_finished = 0
+ total_convs = len(input_conv)
+
+ debug_stats = DebugStats(logger, min(15 * bench_args.num_clients, 500))
+
+ while num_clients_finished < bench_args.num_clients:
+ # Collect updated conversation
+ conv_id, messages = conv_queue.get()
+
+ # Collect results (measurements)
+ while not result_queue.empty():
+ new_data = result_queue.get()
+ client_metrics.append(new_data)
+ debug_stats.update(new_data)
+
+ if conv_id is TERM_SIGNAL:
+ num_clients_finished += 1
+ logger.info(
+ f"{Color.CYAN}{num_clients_finished} out of "
+ f"{bench_args.num_clients} clients finished{Color.RESET}"
+ )
+
+ if bench_args.early_stop and not stop_event.is_set():
+ # Once one client finished, stop all other clients.
+ # there is no reason to continue the benchmark with fewer clients.
+ logger.info(
+ f"{Color.YELLOW}Sending termination signal to clients{Color.RESET}"
+ )
+ stop_event.set()
+ else:
+ output_conv[conv_id] = messages
+
+ finished_convs = len(output_conv)
+ percent = finished_convs / total_convs
+
+ # Tuned to control the print rate (can be changed if required)
+ print_cycle = max(3, int(bench_args.num_clients / 4))
+
+ if finished_convs % print_cycle == 0:
+ runtime_sec = nanosec_to_sec(time.perf_counter_ns() - start_time)
+ logger.info(
+ f"{Color.CYAN}Finished {finished_convs} out of {total_convs} conversations ({percent:.0%}), " # noqa: E501
+ f"{num_clients_finished} out of {bench_args.num_clients} clients finished, collected {len(client_metrics)} measurements, runtime {runtime_sec:.3f} sec{Color.RESET}" # noqa: E501
+ )
+
+ rps: Union[str, float] = round(len(client_metrics) / runtime_sec, 3)
+ if len(client_metrics) < (5 * bench_args.num_clients):
+ # Do not estimate the RPS if the number of samples is very low
+ # (threshold can be tuned if needed)
+ rps = "N/A"
+
+ runtime_left_sec: Union[str, float] = round(
+ (runtime_sec / finished_convs) * (total_convs - finished_convs), 3
+ )
+ if percent < 0.05:
+ # If less than 5% of the conversations were not finished,
+ # the estimation will probably be very inaccurate
+ # (threshold can be tuned if needed).
+ runtime_left_sec = "N/A"
+
+ logger.info(
+ f"{Color.CYAN}Estimated req/sec {rps}, estimated runtime left {runtime_left_sec} sec{Color.RESET}" # noqa: E501
+ )
+ debug_stats.print()
+
+ logger.info(
+ f"{Color.CYAN}All {bench_args.num_clients} clients finished{Color.RESET}"
+ )
+
+ # At this point all the clients finished,
+ # collect results (TTFT, TPOT, etc.) from all the clients.
+ # This needs to happen before calling join on the clients
+ # (result_queue should be emptied).
+ while not result_queue.empty():
+ client_metrics.append(result_queue.get())
+
+ logger.info(f"Collected {len(client_metrics)} samples from all the clients")
+
+ # Wait for all clients to finish
+ for client in clients:
+ logger.info(
+ f"{Color.CYAN}Waiting for client {client.name} "
+ f"(is alive: {client.is_alive()}){Color.RESET}"
+ )
+
+ client.join(timeout=120)
+
+ if client.is_alive():
+ logger.warning(
+ f"{Color.YELLOW}Client {client.name} will be terminated{Color.RESET}"
+ )
+ client.terminate()
+
+ exitcode = client.exitcode
+ if exitcode != 0:
+ logger.error(
+ f"{Color.RED}Client {client.name} exited "
+ f"with exit code {exitcode}{Color.RESET}"
+ )
+
+ logger.info(
+ f"All {bench_args.num_clients} clients exited (successfully "
+ f"finished {len(output_conv)} out of {total_convs} conversations)"
+ )
+
+ # Queues should be closed, required to avoid hang at interpreter shutdown
+ unfinished_tasks = 0
+ while not task_queue.empty():
+ task_queue.get()
+ unfinished_tasks += 1
+
+ if unfinished_tasks > 0:
+ # Can happen if not all tasks (conversations) have finished.
+ # May happen if --max-num-requests was used,
+ # or if an error occurred in one of the clients.
+ logger.debug(f"Discarding {unfinished_tasks} unfinished tasks")
+
+ task_queue.close()
+ task_queue.join_thread()
+
+ result_queue.close()
+ result_queue.join_thread()
+
+ conv_queue.close()
+ conv_queue.join_thread()
+
+ return output_conv, client_metrics
+
+
+def get_filename_with_timestamp(label: str, extension: str) -> str:
+ time_now = datetime.now()
+ timestamp = time_now.strftime("%d-%m-%Y_%H-%M-%S")
+ filename = f"{label}__{timestamp}.{extension}"
+ return filename
+
+
+def process_statistics(
+ client_metrics: list[RequestStats],
+ warmup_percentages: list[float],
+ test_params: dict,
+ verbose: bool,
+ gen_conv_args: Optional[GenConvArgs] = None,
+ excel_output: bool = False,
+) -> None:
+ if len(client_metrics) == 0:
+ logger.info("No samples to process")
+ return
+
+ logger.info(f"Processing {len(client_metrics)} samples...")
+
+ raw_data = pd.DataFrame(client_metrics)
+
+ if verbose:
+ # Calculate the time between user turns in each conversation (in a new column)
+ raw_data = raw_data.sort_values(by=["conversation_id", "start_time_ms"])
+ raw_data["time_between_user_turns_sec"] = raw_data.groupby("conversation_id")[
+ "start_time_ms"
+ ].diff()
+
+ # Convert milliseconds to seconds
+ raw_data["time_between_user_turns_sec"] = (
+ raw_data["time_between_user_turns_sec"] / 1000.0
+ )
+
+ # Final raw data should be sorted by time
+ raw_data = raw_data.sort_values(by=["start_time_ms"])
+ raw_data["end_time_ms"] = raw_data["start_time_ms"] + raw_data["latency_ms"]
+
+ percentiles = [0.25, 0.5, 0.75, 0.9]
+
+ # Add more percentiles if there are enough samples
+ if len(raw_data) >= 100:
+ percentiles.append(0.99)
+
+ if len(raw_data) >= 1000:
+ percentiles.append(0.999)
+
+ if len(raw_data) >= 10000:
+ percentiles.append(0.9999)
+
+ # Set precision for numbers in the output text (the dataframes)
+ pd.set_option("display.precision", 2)
+
+ # Exclude parameters from RequestStats
+ exclude = [
+ "start_time_ms",
+ "end_time_ms",
+ "output_num_first_chunk_tokens",
+ "approx_cached_percent",
+ "conversation_id",
+ "client_id",
+ ]
+
+ print(TEXT_SEPARATOR)
+ print(f"{Color.YELLOW}Parameters:{Color.RESET}")
+ for k, v in test_params.items():
+ print(f"{k}={v}")
+
+ # conversations generation parameters
+ if gen_conv_args is not None:
+ gen_params = {
+ "text_files": ", ".join(gen_conv_args.text_files),
+ "input_num_turns": str(gen_conv_args.input_num_turns),
+ "input_common_prefix_num_tokens": str(
+ gen_conv_args.input_common_prefix_num_tokens
+ ),
+ "input_prefix_num_tokens": str(gen_conv_args.input_prefix_num_tokens),
+ "input_num_tokens": str(gen_conv_args.input_num_tokens),
+ "output_num_tokens": str(gen_conv_args.output_num_tokens),
+ }
+
+ print(f"{Color.YELLOW}Conversations Generation Parameters:{Color.RESET}")
+ for k, v in gen_params.items():
+ print(f"{k}={v}")
+
+ print(TEXT_SEPARATOR)
+
+ params_list = []
+ df_list = []
+ for percent in warmup_percentages:
+ # Select samples from the end (tail) of the dataframe
+ warmup_count = int(percent * len(raw_data))
+ tail_count = len(raw_data) - warmup_count
+ if tail_count == 0:
+ # No reason to process if the count of samples is zero
+ break
+
+ df = raw_data.tail(tail_count)
+
+ # Runtime is the diff between the end of the last request
+ # and the start of the first request
+ runtime_sec = df["end_time_ms"].iloc[-1] - df["start_time_ms"].iloc[0]
+
+ # Convert milliseconds to seconds
+ runtime_sec = runtime_sec / 1000.0
+ requests_per_sec = float(len(df)) / runtime_sec
+
+ params = {"runtime_sec": runtime_sec, "requests_per_sec": requests_per_sec}
+
+ # Generate a summary of relevant metrics (and drop irrelevant data)
+ df = df.drop(columns=exclude).describe(percentiles=percentiles).transpose()
+
+ # List for Excel file
+ params_list.append(params)
+ df_list.append(df)
+
+ # Print the statistics summary
+ if percent > 0 or len(warmup_percentages) > 1:
+ print(
+ f"{Color.YELLOW}Statistics summary "
+ f"(assuming {percent:.0%} warmup samples):{Color.RESET}"
+ )
+ else:
+ print(f"{Color.YELLOW}Statistics summary:{Color.RESET}")
+
+ for k, v in params.items():
+ if isinstance(v, float):
+ print(f"{k} = {v:.3f}")
+ else:
+ print(f"{k} = {v}")
+ print(TEXT_SEPARATOR)
+ print(df)
+ print(TEXT_SEPARATOR)
+
+ if excel_output:
+ prefix = f"statistics_{test_params['num_clients']}_clients"
+ filename = get_filename_with_timestamp(prefix, "xlsx")
+
+ with pd.ExcelWriter(filename, engine="xlsxwriter") as writer:
+ startrow = 0
+ test_params_df = pd.DataFrame([test_params])
+ test_params_df.to_excel(
+ writer, sheet_name="Summary", index=False, startrow=startrow
+ )
+ startrow += len(test_params_df) + 3
+
+ if gen_conv_args is not None:
+ gen_params_df = pd.DataFrame([gen_params])
+ gen_params_df.to_excel(
+ writer, sheet_name="Summary", index=False, startrow=(startrow - 1)
+ )
+ startrow += len(gen_params_df) + 3
+
+ for params, df_stats in zip(params_list, df_list):
+ df_params = pd.DataFrame([params])
+ df_params.to_excel(
+ writer, sheet_name="Summary", index=False, startrow=startrow
+ )
+ startrow += len(df_params) + 2
+ df_stats.to_excel(
+ writer, sheet_name="Summary", index=True, startrow=startrow
+ )
+ startrow += len(df_stats) + 3
+
+ raw_data.to_excel(writer, sheet_name="Raw data", index=False, startrow=0)
+
+ logger.info(
+ f"{Color.GREEN}Client metrics exported to file: {filename}{Color.RESET}"
+ )
+
+
+async def get_server_info(url: str) -> None:
+ logger.info(f"{Color.BLUE}Collecting information from server: {url}{Color.RESET}")
+ async with aiohttp.ClientSession() as session:
+ # Get server version (not mandatory, "version" endpoint may not exist)
+ url_version = f"{url}/version"
+ async with session.get(url_version) as response:
+ if HTTPStatus(response.status) == HTTPStatus.OK:
+ text = await response.text()
+ logger.info(f"{Color.BLUE}Server version: {text}{Color.RESET}")
+
+ # Get available models
+ url_models = f"{url}/v1/models"
+ async with session.get(url_models) as response:
+ if HTTPStatus(response.status) == HTTPStatus.OK:
+ text = await response.text()
+ logger.info(f"{Color.BLUE}Models:{Color.RESET}")
+ models_data = json.loads(text)
+ models_list = models_data["data"]
+ for model in models_list:
+ model_id = model["id"]
+ max_model_len = model.get("max_model_len", "N/A")
+ logger.info(
+ f"{Color.BLUE}\t{model_id=}, {max_model_len=}{Color.RESET}"
+ )
+ else:
+ logger.info(f"{Color.RED}Failed to get models{Color.RESET}")
+
+
+async def main() -> None:
+ parser = argparse.ArgumentParser(
+ prog="Benchmark serving with multi-turn conversations",
+ description="Benchmark online inference using REST API",
+ )
+ parser.add_argument("--version", action="version", version="%(prog)s 1.0")
+
+ parser.add_argument(
+ "-i",
+ "--input-file",
+ type=str,
+ required=True,
+ help="Input JSON file with ShareGPT conversations or "
+ "configuration file for generation of synthetic conversations",
+ )
+ parser.add_argument(
+ "-o",
+ "--output-file",
+ type=str,
+ default=None,
+ help="Output JSON file containing conversations with updated assistant answers",
+ )
+
+ parser.add_argument(
+ "--seed",
+ type=int,
+ default=0,
+ help="Seed for random number generators (default: 0)",
+ )
+
+ parser.add_argument(
+ "-m", "--model", type=str, required=True, help="Path of the LLM model"
+ )
+ parser.add_argument(
+ "--served-model-name",
+ type=str,
+ default=None,
+ help="The model name used in the API. "
+ "If not specified, the model name will be the "
+ "same as the ``--model`` argument. ",
+ )
+
+ parser.add_argument(
+ "-u",
+ "--url",
+ type=str,
+ default="http://localhost:8000",
+ help="Base URL for the LLM API server",
+ )
+
+ parser.add_argument(
+ "-p",
+ "--num-clients",
+ type=int,
+ default=1,
+ help="Number of clients that will send requests in parallel",
+ )
+ parser.add_argument(
+ "-k",
+ "--max-active-conversations",
+ type=int,
+ default=None,
+ help="Max number of active conversations at a time (for all clients)",
+ )
+ parser.add_argument(
+ "-n",
+ "--max-num-requests",
+ type=int,
+ default=None,
+ help="Max number of requests to send (total for all clients)",
+ )
+
+ parser.add_argument(
+ "--warmup-step",
+ default=False,
+ action="store_true",
+ help="Run a warmup step (using only the first turn of every conversation), "
+ "measurements will not be included in the final benchmark results",
+ )
+
+ parser.add_argument(
+ "--max-turns",
+ type=int,
+ default=None,
+ help="Maximum number of turns/messages per conversation, "
+ "includes both user and assistant messages "
+ "(a positive number, e.g: 2, 4, 6, etc.), disabled by default",
+ )
+ parser.add_argument(
+ "--no-early-stop",
+ default=False,
+ action="store_true",
+ help="By default, the benchmark will stop if at least one client exits."
+ " Use this flag to disable this behavior",
+ )
+
+ parser.add_argument(
+ "--limit-max-tokens",
+ type=int,
+ default=NUM_TOKENS_FROM_DATASET,
+ help="Set max_tokens for the output token count of each request "
+ "(must also set --limit-min-tokens). "
+ "Overrides output token count from the input dataset. "
+ "Use a negative value to disable this limit.",
+ )
+ parser.add_argument(
+ "--limit-min-tokens",
+ type=int,
+ default=NUM_TOKENS_FROM_DATASET,
+ help="Set min_tokens for the output token count of each request "
+ "(must also set --limit-max-tokens). "
+ "Overrides output token count from the input dataset. "
+ "Use a negative value to disable this limit.",
+ )
+
+ parser.add_argument(
+ "--request-rate",
+ type=float,
+ default=0,
+ help="Expected request rate (Poisson process) per client in requests/sec."
+ "Set to 0 for no delay between requests.",
+ )
+ parser.add_argument(
+ "--conversation-sampling",
+ type=ConversationSampling,
+ choices=list(ConversationSampling),
+ default=ConversationSampling.ROUND_ROBIN,
+ help=(
+ "Strategy for selecting which conversation to use for the next request. "
+ "Options: 'round_robin' (cycle through conversations), "
+ "'random' (pick randomly)."
+ ),
+ )
+ parser.add_argument(
+ "--verify-output",
+ default=False,
+ action="store_true",
+ help="Verify the LLM output (compare to the answers in the input JSON file)",
+ )
+
+ parser.add_argument(
+ "--no-stream",
+ default=False,
+ action="store_true",
+ help="Disable stream/streaming mode (set 'stream' to False in the API request)",
+ )
+
+ parser.add_argument(
+ "-e",
+ "--excel-output",
+ default=False,
+ action="store_true",
+ help="Export summary to Excel file (optional)",
+ )
+ parser.add_argument(
+ "-v",
+ "--verbose",
+ default=False,
+ action="store_true",
+ help="Enable verbose output",
+ )
+ parser.add_argument(
+ "--print-content",
+ default=False,
+ action="store_true",
+ help="Print the user prompts and the server's answers",
+ )
+
+ parser.add_argument(
+ "--warmup-percentages",
+ type=str,
+ default="0%",
+ help="Ignore the first X samples as warmup (X is a percentage)."
+ " A comma separated list of percentages can be used "
+ "(for example: --warmup-percentages=0%%,50%%)",
+ )
+
+ args = parser.parse_args()
+
+ logger.info(args)
+
+ logger.info(f"{Color.GREEN}Input parameters:{Color.RESET}")
+ logger.info(f"url={args.url}")
+ logger.info(f"model={args.model}")
+ logger.info(f"num_clients={args.num_clients}")
+
+ if args.verify_output:
+ logger.info(f"{Color.PURPLE}Verify is enabled{Color.RESET}")
+
+ # Calculate the amount of samples to filter (as warmup samples/measurements).
+ try:
+ warmup_percentages: list[float] = [0.0]
+ if not args.warmup_step:
+ # Warmup percentage can be used only if the warmup step was used
+ warmup_strings: list[str] = args.warmup_percentages.split(",")
+ warmup_strings = [x.replace("%", "") for x in warmup_strings]
+ warmup_percentages = [float(x) / 100 for x in warmup_strings]
+
+ # Check for valid range (0 to 1)
+ for p in warmup_percentages:
+ assert p >= 0.0 and p < 1.0
+
+ # Sort from high to low warmup percentage
+ warmup_percentages.sort()
+
+ logger.info(
+ f"Warmup percentages (percentage of samples): {warmup_percentages}"
+ )
+
+ except Exception:
+ raise ValueError(
+ f"Invalid --warmup-percentage={args.warmup_percentage}"
+ ) from None
+
+ random.seed(args.seed)
+ np.random.seed(args.seed)
+
+ if not os.path.exists(args.model):
+ raise OSError(f"Path does not exist: {args.model}")
+ logger.info("Loading tokenizer")
+ tokenizer = AutoTokenizer.from_pretrained(args.model)
+
+ await get_server_info(args.url)
+
+ # Load the input file (either conversations of configuration file)
+ logger.info(f"Reading input file: {args.input_file}")
+ with open(args.input_file) as f:
+ input_data = json.load(f)
+
+ gen_conv_args = None
+ if isinstance(input_data, list):
+ # The conversations are stored as a list of dicts
+ logger.info(f"Found {len(input_data)} items in the input file")
+
+ # Convert the list to a ConversationsMap
+ conversations = conversations_list_to_dict(input_data)
+
+ elif isinstance(input_data, dict):
+ # The input file is a configuration file
+ # (type is determined by the field 'filetype')
+ if "filetype" not in input_data:
+ raise Exception(
+ f"Input file {args.input_file} is invalid (missing 'filetype')"
+ )
+
+ logger.info(f"Using input file with filetype: {input_data['filetype']}")
+
+ gen_conv_args = parse_input_json_file(input_data)
+
+ # Disable warning from "huggingface/tokenizers"
+ # (when using python multiprocessing and tokenizers)
+ os.environ["TOKENIZERS_PARALLELISM"] = "true"
+
+ # Generate synthetic conversations
+ conversations = generate_conversations(gen_conv_args, tokenizer)
+
+ else:
+ raise Exception(f"Input file {args.input_file} is invalid")
+
+ if args.max_turns is not None:
+ if args.max_turns < 1:
+ raise ValueError("Max turns must be a positive number")
+ logger.info(
+ f"{Color.PURPLE}Max turns per conversation "
+ f"is limited to {args.max_turns}{Color.RESET}"
+ )
+
+ # Create benchmark configurations
+ client_args, req_args = get_client_config(args, conversations)
+
+ bench_args = BenchmarkArgs(
+ url=args.url, num_clients=args.num_clients, early_stop=not args.no_early_stop
+ )
+
+ # Warm-up step
+ if args.warmup_step:
+ # Only send a single user prompt from every conversation.
+ # max_active_conversations must be 1,
+ # otherwise the clients may exit after sending a single request
+ # (because the task queue is empty).
+ warmup_client_args = client_args._replace(
+ skip_first_turn=False, max_turns=1, max_active_conversations=1
+ )
+
+ # Early stop should be disabled,
+ # all clients should finish their work before exiting
+ warmup_bench_args = bench_args._replace(early_stop=False)
+
+ logger.info(f"{Color.PURPLE}Warmup start{Color.RESET}")
+ conversations, _ = await main_mp(
+ warmup_client_args, req_args, warmup_bench_args, tokenizer, conversations
+ )
+ logger.info(f"{Color.PURPLE}Warmup done{Color.RESET}")
+
+ # Run the benchmark
+ start_time = time.perf_counter_ns()
+ client_convs, client_metrics = await main_mp(
+ client_args, req_args, bench_args, tokenizer, conversations
+ )
+ total_runtime_ms = nanosec_to_millisec(time.perf_counter_ns() - start_time)
+
+ # Calculate requests per second
+ total_runtime_sec = total_runtime_ms / 1000.0
+ rps = len(client_metrics) / total_runtime_sec
+ logger.info(
+ f"{Color.GREEN}All clients finished, total runtime: {total_runtime_sec:.3f} sec"
+ f" ({total_runtime_ms:.3f} ms), requests per second: {rps:.3f}{Color.RESET}"
+ )
+
+ # Benchmark parameters
+ params = {
+ "model": args.model,
+ "num_clients": args.num_clients,
+ "num_conversations": len(conversations),
+ "active_conversations": args.max_active_conversations,
+ "seed": args.seed,
+ }
+
+ if args.limit_min_tokens > 0:
+ params["min_tokens"] = args.limit_min_tokens
+
+ if args.limit_max_tokens > 0:
+ params["max_tokens"] = args.limit_max_tokens
+
+ # Process and print statistics (and save excel file with the statistics)
+ process_statistics(
+ client_metrics,
+ test_params=params,
+ warmup_percentages=warmup_percentages,
+ verbose=args.verbose,
+ gen_conv_args=gen_conv_args,
+ excel_output=args.excel_output,
+ )
+
+ if args.output_file is not None:
+ # Write a JSON file with the updated conversations
+ # The "assistant" content will contain the answers from the tested LLM
+ output_data: ShareGptConversations = conversations_dict_to_list(client_convs)
+ logger.info(
+ f"{Color.GREEN}Writing conversations file: {args.output_file}{Color.RESET}"
+ )
+ with open(args.output_file, "w") as f:
+ json.dump(output_data, f, indent=4)
+
+
+if __name__ == "__main__":
+ asyncio.run(main())
diff --git a/benchmarks/multi_turn/convert_sharegpt_to_openai.py b/benchmarks/multi_turn/convert_sharegpt_to_openai.py
new file mode 100644
index 0000000000..c3622c99a2
--- /dev/null
+++ b/benchmarks/multi_turn/convert_sharegpt_to_openai.py
@@ -0,0 +1,354 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+Download dataset from:
+https://huggingface.co/datasets/philschmid/sharegpt-raw/blob/main/sharegpt_20230401_clean_lang_split.json
+
+Convert to OpenAI API:
+export INPUT_FILE=sharegpt_20230401_clean_lang_split.json
+python convert_sharegpt_to_openai.py $INPUT_FILE sharegpt_conv_128.json --max-items=128
+"""
+
+import argparse
+import json
+import random
+from statistics import mean
+from typing import Any, Optional
+
+import pandas as pd # type: ignore
+import tqdm # type: ignore
+from transformers import AutoTokenizer # type: ignore
+
+
+def has_non_english_chars(text: str) -> bool:
+ return not text.isascii()
+
+
+def content_is_valid(
+ content: str, min_content_len: Optional[int], max_content_len: Optional[int]
+) -> bool:
+ if min_content_len and len(content) < min_content_len:
+ return False
+
+ if max_content_len and len(content) > max_content_len:
+ return False
+
+ return has_non_english_chars(content)
+
+
+def print_stats(
+ conversations: "list[dict[Any, Any]]", tokenizer: Optional[AutoTokenizer] = None
+) -> None:
+ # Collect statistics
+ stats = []
+
+ print("\nCollecting statistics...")
+ for item in tqdm.tqdm(conversations):
+ # item has "id" and "messages"
+ messages = item["messages"]
+
+ user_turns = 0
+ assistant_turns = 0
+ user_words = 0
+ assistant_words = 0
+ conv_chars = 0
+
+ user_tokens: list[int] = []
+ assistant_tokens: list[int] = []
+
+ for m in messages:
+ content = m["content"]
+ conv_chars += len(content)
+ content_num_words = content.count(" ") + 1
+
+ num_tokens = 0
+ if tokenizer:
+ num_tokens = len(tokenizer(m["content"]).input_ids)
+
+ if m["role"] == "user":
+ user_turns += 1
+ user_words += content_num_words
+ if tokenizer:
+ user_tokens.append(num_tokens)
+
+ elif m["role"] == "assistant":
+ assistant_turns += 1
+ assistant_words += content_num_words
+ if tokenizer:
+ assistant_tokens.append(num_tokens)
+
+ # assert user_turns == assistant_turns, \
+ # f"Invalid conversation ID {item['id']}"
+
+ conv_words = user_words + assistant_words
+ item_stats = {
+ "user_turns": user_turns,
+ "assistant_turns": assistant_turns,
+ "user_words": user_words,
+ "assistant_words": assistant_words,
+ "conv_turns": len(messages),
+ "conv_words": conv_words,
+ "conv_characters": conv_chars,
+ }
+
+ if len(user_tokens) > 0:
+ item_stats["user_tokens"] = int(mean(user_tokens))
+
+ if len(assistant_tokens) > 0:
+ item_stats["assistant_tokens"] = int(mean(assistant_tokens))
+
+ stats.append(item_stats)
+
+ print("\nStatistics:")
+ percentiles = [0.25, 0.5, 0.75, 0.9, 0.99, 0.999, 0.9999]
+ df = pd.DataFrame(stats)
+ print(df.describe(percentiles=percentiles).transpose())
+
+
+def convert_sharegpt_to_openai(
+ seed: int,
+ input_file: str,
+ output_file: str,
+ max_items: Optional[int],
+ min_content_len: Optional[int] = None,
+ max_content_len: Optional[int] = None,
+ min_turns: Optional[int] = None,
+ max_turns: Optional[int] = None,
+ model: Optional[str] = None,
+) -> None:
+ if min_turns and max_turns:
+ assert min_turns <= max_turns
+
+ if min_content_len and max_content_len:
+ # Verify that min is not larger than max if both were given
+ assert min_content_len <= max_content_len
+
+ print(
+ f"Input parameters:\n{seed=}, {max_items=}, {min_content_len=},"
+ f" {max_content_len=}, {min_turns=}, {max_turns=}\n"
+ )
+
+ random.seed(seed)
+
+ tokenizer = None
+ if model is not None:
+ print(f"Loading tokenizer from: {model}")
+ tokenizer = AutoTokenizer.from_pretrained(model)
+
+ # Read the ShareGPT JSON file
+ print(f"Reading file: {input_file}")
+ with open(input_file, encoding="utf-8") as f:
+ # Should be a list of dicts
+ # Each dict should have "id" (string) and "conversations" (list of dicts)
+ sharegpt_data = json.load(f)
+
+ assert isinstance(sharegpt_data, list), "Input file should contain a list of dicts"
+
+ print(f"Total items in input file: {len(sharegpt_data):,}")
+
+ print(f"Shuffling dataset with seed {seed}")
+ random.shuffle(sharegpt_data)
+
+ # Map conversation ID to the all the messages
+ conversation_parts: dict[str, list[Any]] = {}
+
+ for item in tqdm.tqdm(sharegpt_data):
+ assert "id" in item, "Missing key 'id'"
+ assert "conversations" in item, "Missing key 'conversations'"
+
+ # Conversation ID (e.g: "hiWPlMD") and part/session (0, 1, 2, etc.)
+ conv_id, _ = item["id"].split("_")
+ new_turns = item["conversations"]
+
+ if conv_id not in conversation_parts:
+ # Start new conversation
+ conversation_parts[conv_id] = []
+ elif len(conversation_parts[conv_id]) > 0 and len(new_turns) > 0:
+ prev_turns = conversation_parts[conv_id][-1]
+ if prev_turns[-1]["from"] == new_turns[0]["from"]:
+ new_turns = new_turns[1:]
+
+ if len(new_turns) > 0:
+ # We assume that parts are in order in the ShareGPT dataset
+ conversation_parts[conv_id].append(new_turns)
+
+ dataset: list[dict[str, Any]] = []
+ for conv_id, conv_parts in conversation_parts.items():
+ new_item = {"id": conv_id}
+
+ conversations: list[dict[str, str]] = []
+
+ # Merge all parts
+ for conv_part in conv_parts:
+ conversations.extend(conv_part)
+
+ if len(conversations) > 0:
+ new_item["conversations"] = conversations
+ dataset.append(new_item)
+
+ print(f"Total unique conversations (IDs) in input file: {len(dataset):,}")
+
+ # Final output data
+ final_openai_dataset: list[dict] = []
+
+ # Filter conversations from the ShareGPT dataset and convert to OpenAI format
+ for item in tqdm.tqdm(dataset):
+ messages: list[dict] = []
+
+ assert "id" in item, "Missing key 'id'"
+ assert "conversations" in item, "Missing key 'conversations'"
+
+ conv_id = item["id"]
+ conversations = item["conversations"]
+
+ if min_turns is not None and len(conversations) < min_turns:
+ # Skip short conversations
+ continue
+
+ # Convert each message in the conversation, up to max_turns if specified
+ for i, turn in enumerate(conversations):
+ assert "from" in turn and "value" in turn, (
+ f"Invalid conversation ID {conv_id} - missing 'from' or 'value'"
+ )
+
+ role = None
+ turn_from = turn["from"]
+
+ if turn_from in {"human", "user"}:
+ role = "user"
+ elif turn_from in {"gpt", "bing", "chatgpt", "bard"}:
+ role = "assistant"
+ elif turn_from == "system":
+ role = "system"
+
+ assert role is not None, (
+ f"Invalid conversation ID {conv_id} - 'from'='{turn_from}' is invalid"
+ )
+
+ if i == 0 and role != "user":
+ # If the first message is from assistant (gpt), skip it.
+ # this happens when the conversation is a follow-up
+ # to a previous conversation (from the same user).
+ continue
+
+ if max_turns is not None and i >= max_turns:
+ break
+
+ # Convert message to OpenAI format (with "role" and "content")
+ content = turn["value"]
+ messages.append({"role": role, "content": content})
+
+ # Add the converted conversation to the OpenAI format
+ if len(messages) > 0:
+ valid_messages = True
+
+ # First turn should always be from the user
+ user_turn = True
+
+ for m in messages:
+ # Make sure that turns alternate between user and assistant
+ if (user_turn and m["role"] != "user") or (
+ not user_turn and m["role"] != "assistant"
+ ):
+ valid_messages = False
+ break
+
+ user_turn = not user_turn
+
+ content = m["content"]
+ valid_messages = content_is_valid(
+ content, min_content_len, max_content_len
+ )
+ if not valid_messages:
+ break
+
+ if valid_messages is True:
+ final_openai_dataset.append({"id": conv_id, "messages": messages})
+
+ assert len(final_openai_dataset) > 0, "Final number of conversations is zero"
+
+ print_stats(final_openai_dataset)
+
+ print_stats_again = False
+ if max_items is not None and len(final_openai_dataset) > max_items:
+ print(f"\n\nSampling {max_items} items from the dataset...")
+ print_stats_again = True
+ final_openai_dataset = random.sample(final_openai_dataset, max_items)
+
+ if print_stats_again:
+ # Print stats after the dataset changed
+ print_stats(final_openai_dataset, tokenizer)
+
+ # Write the converted data to a new JSON file
+ final_size = len(final_openai_dataset)
+ print(f"\nTotal conversations converted (after filtering): {final_size:,}")
+ print(f"\nWriting file: {output_file}")
+ with open(output_file, "w", encoding="utf-8") as f:
+ json.dump(final_openai_dataset, f, ensure_ascii=False, indent=2)
+
+
+def main() -> None:
+ parser = argparse.ArgumentParser(
+ description="Convert ShareGPT dataset to OpenAI API format"
+ )
+ parser.add_argument("input_file", help="Path to the input ShareGPT JSON file")
+ parser.add_argument(
+ "output_file", help="Path to the output OpenAI format JSON file"
+ )
+ parser.add_argument(
+ "--seed", type=int, default=0, help="Seed for random number generators"
+ )
+ parser.add_argument(
+ "--max-items",
+ type=int,
+ default=None,
+ help="Maximum number of items in the output file",
+ )
+ parser.add_argument(
+ "--min-turns",
+ type=int,
+ default=None,
+ help="Minimum number of turns per conversation",
+ )
+ parser.add_argument(
+ "--max-turns",
+ type=int,
+ default=None,
+ help="Maximum number of turns per conversation",
+ )
+ parser.add_argument(
+ "--min-content-len",
+ type=int,
+ default=None,
+ help="Min number of characters in the messages' content",
+ )
+ parser.add_argument(
+ "--max-content-len",
+ type=int,
+ default=None,
+ help="Max number of characters in the messages' content",
+ )
+ parser.add_argument(
+ "--model",
+ type=str,
+ default=None,
+ help="LLM model, only the tokenizer will be used",
+ )
+
+ args = parser.parse_args()
+
+ convert_sharegpt_to_openai(
+ args.seed,
+ args.input_file,
+ args.output_file,
+ args.max_items,
+ args.min_content_len,
+ args.max_content_len,
+ args.min_turns,
+ args.max_turns,
+ args.model,
+ )
+
+
+if __name__ == "__main__":
+ main()
diff --git a/benchmarks/multi_turn/generate_multi_turn.json b/benchmarks/multi_turn/generate_multi_turn.json
new file mode 100644
index 0000000000..274d03c2bd
--- /dev/null
+++ b/benchmarks/multi_turn/generate_multi_turn.json
@@ -0,0 +1,35 @@
+{
+ "filetype": "generate_conversations",
+ "num_conversations": 24,
+ "text_files": ["pg1184.txt"],
+ "print_stats": false,
+ "prompt_input": {
+ "num_turns": {
+ "distribution": "uniform",
+ "min": 12,
+ "max": 18
+ },
+ "common_prefix_num_tokens": {
+ "distribution": "constant",
+ "value": 500
+ },
+ "prefix_num_tokens": {
+ "distribution": "lognormal",
+ "mean": 6,
+ "sigma": 4,
+ "max": 1500
+ },
+ "num_tokens": {
+ "distribution": "uniform",
+ "min": 120,
+ "max": 160
+ }
+ },
+ "prompt_output": {
+ "num_tokens": {
+ "distribution": "uniform",
+ "min": 80,
+ "max": 120
+ }
+ }
+}
\ No newline at end of file
diff --git a/benchmarks/multi_turn/requirements.txt b/benchmarks/multi_turn/requirements.txt
new file mode 100644
index 0000000000..f0e1935914
--- /dev/null
+++ b/benchmarks/multi_turn/requirements.txt
@@ -0,0 +1,5 @@
+numpy>=1.24
+pandas>=2.0.0
+aiohttp>=3.10
+transformers>=4.46
+xlsxwriter>=3.2.1
\ No newline at end of file
diff --git a/cmake/cpu_extension.cmake b/cmake/cpu_extension.cmake
index e0da46e2ac..0649446322 100644
--- a/cmake/cpu_extension.cmake
+++ b/cmake/cpu_extension.cmake
@@ -1,6 +1,7 @@
include(FetchContent)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CXX_EXTENSIONS ON)
set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
@@ -87,6 +88,7 @@ is_avx512_disabled(AVX512_DISABLED)
if (MACOSX_FOUND AND CMAKE_SYSTEM_PROCESSOR STREQUAL "arm64")
message(STATUS "Apple Silicon Detected")
+ set(APPLE_SILICON_FOUND TRUE)
set(ENABLE_NUMA OFF)
check_sysctl(hw.optional.neon ASIMD_FOUND)
check_sysctl(hw.optional.arm.FEAT_BF16 ARM_BF16_FOUND)
@@ -182,17 +184,17 @@ endif()
#
# Build oneDNN for W8A8 GEMM kernels (only for x86-AVX512 /ARM platforms)
# Flag to enable ACL kernels for AARCH64 platforms
-if ( VLLM_BUILD_ACL STREQUAL "ON")
+if (VLLM_BUILD_ACL STREQUAL "ON")
set(USE_ACL ON)
else()
set(USE_ACL OFF)
endif()
-if ((AVX512_FOUND AND NOT AVX512_DISABLED) OR ASIMD_FOUND)
+if ((AVX512_FOUND AND NOT AVX512_DISABLED) OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND OR POWER10_FOUND OR POWER11_FOUND)
FetchContent_Declare(
oneDNN
GIT_REPOSITORY https://github.com/oneapi-src/oneDNN.git
- GIT_TAG v3.8.1
+ GIT_TAG v3.9
GIT_PROGRESS TRUE
GIT_SHALLOW TRUE
)
@@ -204,7 +206,7 @@ if ((AVX512_FOUND AND NOT AVX512_DISABLED) OR ASIMD_FOUND)
endif()
set(ONEDNN_AARCH64_USE_ACL "ON")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wl,-rpath,$ENV{ACL_ROOT_DIR}/build/")
- endif()
+ endif()
set(ONEDNN_LIBRARY_TYPE "STATIC")
set(ONEDNN_BUILD_DOC "OFF")
@@ -217,38 +219,23 @@ if ((AVX512_FOUND AND NOT AVX512_DISABLED) OR ASIMD_FOUND)
set(ONEDNN_ENABLE_ITT_TASKS "OFF")
set(ONEDNN_ENABLE_MAX_CPU_ISA "OFF")
set(ONEDNN_ENABLE_CPU_ISA_HINTS "OFF")
+ set(ONEDNN_VERBOSE "OFF")
set(CMAKE_POLICY_DEFAULT_CMP0077 NEW)
FetchContent_MakeAvailable(oneDNN)
-
- list(APPEND LIBS dnnl)
-elseif(POWER10_FOUND)
- FetchContent_Declare(
- oneDNN
- GIT_REPOSITORY https://github.com/oneapi-src/oneDNN.git
- GIT_TAG v3.7.2
- GIT_PROGRESS TRUE
- GIT_SHALLOW TRUE
+ add_library(dnnl_ext OBJECT "csrc/cpu/dnnl_helper.cpp")
+ target_include_directories(
+ dnnl_ext
+ PUBLIC ${oneDNN_SOURCE_DIR}/include
+ PUBLIC ${oneDNN_BINARY_DIR}/include
+ PRIVATE ${oneDNN_SOURCE_DIR}/src
)
-
- set(ONEDNN_LIBRARY_TYPE "STATIC")
- set(ONEDNN_BUILD_DOC "OFF")
- set(ONEDNN_BUILD_EXAMPLES "OFF")
- set(ONEDNN_BUILD_TESTS "OFF")
- set(ONEDNN_ENABLE_WORKLOAD "INFERENCE")
- set(ONEDNN_ENABLE_PRIMITIVE "MATMUL;REORDER")
- set(ONEDNN_BUILD_GRAPH "OFF")
- set(ONEDNN_ENABLE_JIT_PROFILING "OFF")
- set(ONEDNN_ENABLE_ITT_TASKS "OFF")
- set(ONEDNN_ENABLE_MAX_CPU_ISA "OFF")
- set(ONEDNN_ENABLE_CPU_ISA_HINTS "OFF")
- set(CMAKE_POLICY_DEFAULT_CMP0077 NEW)
-
- set(DNNL_CPU_RUNTIME "OMP")
-
- FetchContent_MakeAvailable(oneDNN)
-
- list(APPEND LIBS dnnl)
+ target_link_libraries(dnnl_ext dnnl)
+ target_compile_options(dnnl_ext PRIVATE ${CXX_COMPILE_FLAGS} -fPIC)
+ list(APPEND LIBS dnnl_ext)
+ set(USE_ONEDNN ON)
+else()
+ set(USE_ONEDNN OFF)
endif()
message(STATUS "CPU extension compile flags: ${CXX_COMPILE_FLAGS}")
@@ -275,7 +262,6 @@ set(VLLM_EXT_SRC
if (AVX512_FOUND AND NOT AVX512_DISABLED)
set(VLLM_EXT_SRC
- "csrc/cpu/quant.cpp"
"csrc/cpu/shm.cpp"
${VLLM_EXT_SRC})
if (ENABLE_AVX512BF16 AND ENABLE_AVX512VNNI)
@@ -289,14 +275,11 @@ if (AVX512_FOUND AND NOT AVX512_DISABLED)
${VLLM_EXT_SRC})
add_compile_definitions(-DCPU_CAPABILITY_AVX512)
endif()
-elseif(POWER10_FOUND)
- set(VLLM_EXT_SRC
- "csrc/cpu/quant.cpp"
- ${VLLM_EXT_SRC})
endif()
-if (ASIMD_FOUND)
+
+if(USE_ONEDNN)
set(VLLM_EXT_SRC
- "csrc/cpu/quant.cpp"
+ "csrc/cpu/dnnl_kernels.cpp"
${VLLM_EXT_SRC})
endif()
diff --git a/cmake/external_projects/flashmla.cmake b/cmake/external_projects/flashmla.cmake
index 6291475164..02224cfe3e 100644
--- a/cmake/external_projects/flashmla.cmake
+++ b/cmake/external_projects/flashmla.cmake
@@ -19,7 +19,7 @@ else()
FetchContent_Declare(
flashmla
GIT_REPOSITORY https://github.com/vllm-project/FlashMLA.git
- GIT_TAG 575f7724b9762f265bbee5889df9c7d630801845
+ GIT_TAG a757314c04eedd166e329e846c820eb1bdd702de
GIT_PROGRESS TRUE
CONFIGURE_COMMAND ""
BUILD_COMMAND ""
@@ -37,13 +37,14 @@ cuda_archs_loose_intersection(FLASH_MLA_ARCHS "9.0a" "${CUDA_ARCHS}")
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER 12.3 AND FLASH_MLA_ARCHS)
set(FlashMLA_SOURCES
${flashmla_SOURCE_DIR}/csrc/flash_api.cpp
- ${flashmla_SOURCE_DIR}/csrc/flash_fwd_mla_bf16_sm90.cu
- ${flashmla_SOURCE_DIR}/csrc/flash_fwd_mla_fp16_sm90.cu
- ${flashmla_SOURCE_DIR}/csrc/flash_fwd_mla_metadata.cu)
+ ${flashmla_SOURCE_DIR}/csrc/kernels/get_mla_metadata.cu
+ ${flashmla_SOURCE_DIR}/csrc/kernels/mla_combine.cu
+ ${flashmla_SOURCE_DIR}/csrc/kernels/splitkv_mla.cu
+ ${flashmla_SOURCE_DIR}/csrc/kernels_fp8/flash_fwd_mla_fp8_sm90.cu)
set(FlashMLA_INCLUDES
${flashmla_SOURCE_DIR}/csrc/cutlass/include
- ${flashmla_SOURCE_DIR}/csrc/include)
+ ${flashmla_SOURCE_DIR}/csrc)
set_gencode_flags_for_srcs(
SRCS "${FlashMLA_SOURCES}"
diff --git a/cmake/external_projects/vllm_flash_attn.cmake b/cmake/external_projects/vllm_flash_attn.cmake
index ef45a5fbeb..3d32121f13 100644
--- a/cmake/external_projects/vllm_flash_attn.cmake
+++ b/cmake/external_projects/vllm_flash_attn.cmake
@@ -38,7 +38,7 @@ else()
FetchContent_Declare(
vllm-flash-attn
GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
- GIT_TAG 1c2624e53c078854e0637ee566c72fe2107e75f4
+ GIT_TAG ee4d25bd84e0cbc7e0b9b9685085fd5db2dcb62a
GIT_PROGRESS TRUE
# Don't share the vllm-flash-attn build between build types
BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
diff --git a/cmake/utils.cmake b/cmake/utils.cmake
index 621179a701..9c0ed1d095 100644
--- a/cmake/utils.cmake
+++ b/cmake/utils.cmake
@@ -467,6 +467,12 @@ function (define_gpu_extension_target GPU_MOD_NAME)
if (GPU_LANGUAGE STREQUAL "HIP")
# Make this target dependent on the hipify preprocessor step.
add_dependencies(${GPU_MOD_NAME} hipify${GPU_MOD_NAME})
+ # Make sure we include the hipified versions of the headers, and avoid conflicts with the ones in the original source folder
+ target_include_directories(${GPU_MOD_NAME} PRIVATE ${CMAKE_CURRENT_BINARY_DIR}/csrc
+ ${GPU_INCLUDE_DIRECTORIES})
+ else()
+ target_include_directories(${GPU_MOD_NAME} PRIVATE csrc
+ ${GPU_INCLUDE_DIRECTORIES})
endif()
if (GPU_ARCHITECTURES)
@@ -482,8 +488,6 @@ function (define_gpu_extension_target GPU_MOD_NAME)
target_compile_definitions(${GPU_MOD_NAME} PRIVATE
"-DTORCH_EXTENSION_NAME=${GPU_MOD_NAME}")
- target_include_directories(${GPU_MOD_NAME} PRIVATE csrc
- ${GPU_INCLUDE_DIRECTORIES})
target_link_libraries(${GPU_MOD_NAME} PRIVATE torch ${GPU_LIBRARIES})
diff --git a/csrc/activation_kernels.cu b/csrc/activation_kernels.cu
index 55e6596797..a4a880f13c 100644
--- a/csrc/activation_kernels.cu
+++ b/csrc/activation_kernels.cu
@@ -128,6 +128,45 @@ __global__ void act_and_mul_kernel_with_param(
}
}
+template
+__device__ __forceinline__ T swigluoai_and_mul(const T& gate, const T& up,
+ float alpha, float limit) {
+ // clamp gate: min=None, max=limit
+ const float gate_f = (float)gate;
+ const float clamped_gate = gate_f > limit ? limit : gate_f;
+
+ // clamp up: min=-limit, max=limit
+ const float up_f = (float)up;
+ const float clamped_up =
+ up_f > limit ? limit : (up_f < -limit ? -limit : up_f);
+
+ // glu = gate * sigmoid(gate * alpha)
+ const float sigmoid_val = 1.0f / (1.0f + expf(-clamped_gate * alpha));
+ const float glu = clamped_gate * sigmoid_val;
+
+ // (up + 1) * glu
+ return (T)((clamped_up + 1.0f) * glu);
+}
+
+template
+__global__ void swigluoai_and_mul_kernel(
+ scalar_t* __restrict__ out, // [..., d]
+ const scalar_t* __restrict__ input, // [..., 2, d]
+ const int d, const float alpha, const float limit) {
+ const int64_t token_idx = blockIdx.x;
+ // TODO: Vectorize loads and stores.
+ for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
+ // gate = x[..., ::2] (even indices)
+ const scalar_t gate = VLLM_LDG(&input[token_idx * 2 * d + 2 * idx]);
+ // up = x[..., 1::2] (odd indices)
+ const scalar_t up = VLLM_LDG(&input[token_idx * 2 * d + 2 * idx + 1]);
+
+ out[token_idx * d + idx] = ACT_FN(gate, up, alpha, limit);
+ }
+}
+
} // namespace vllm
#define LAUNCH_ACTIVATION_GATE_KERNEL_WITH_PARAM(KERNEL, PARAM) \
@@ -145,11 +184,31 @@ __global__ void act_and_mul_kernel_with_param(
PARAM); \
});
+#define LAUNCH_SIGLUOAI_AND_MUL(KERNEL, ALPHA, LIMIT) \
+ int d = input.size(-1) / 2; \
+ int64_t num_tokens = input.numel() / input.size(-1); \
+ dim3 grid(num_tokens); \
+ dim3 block(std::min(d, 1024)); \
+ const at::cuda::OptionalCUDAGuard device_guard(device_of(input)); \
+ const cudaStream_t stream = at::cuda::getCurrentCUDAStream(); \
+ VLLM_DISPATCH_FLOATING_TYPES( \
+ input.scalar_type(), "clamp_swiglu_kernel_with_params", [&] { \
+ vllm::swigluoai_and_mul_kernel> \
+ <<>>(out.data_ptr(), \
+ input.data_ptr(), d, ALPHA, \
+ LIMIT); \
+ });
+
void fatrelu_and_mul(torch::Tensor& out, // [..., d],
torch::Tensor& input, // [..., 2 * d]
double threshold) {
LAUNCH_ACTIVATION_GATE_KERNEL_WITH_PARAM(vllm::fatrelu_kernel, threshold);
}
+void swigluoai_and_mul(torch::Tensor& out, // [..., d]
+ torch::Tensor& input, // [..., 2 * d]
+ double alpha, double limit) {
+ LAUNCH_SIGLUOAI_AND_MUL(vllm::swigluoai_and_mul, alpha, limit);
+}
namespace vllm {
// Element-wise activation kernel template.
diff --git a/csrc/attention/mla/sm100_cutlass_mla_kernel.cu b/csrc/attention/mla/sm100_cutlass_mla_kernel.cu
index e0e95d0629..c60f1823b8 100644
--- a/csrc/attention/mla/sm100_cutlass_mla_kernel.cu
+++ b/csrc/attention/mla/sm100_cutlass_mla_kernel.cu
@@ -36,6 +36,7 @@ limitations under the License.
#if !defined(CUDA_VERSION) || CUDA_VERSION < 12040
void sm100_cutlass_mla_decode(
torch::Tensor const& out,
+ torch::Tensor const& lse,
torch::Tensor const& q_nope,
torch::Tensor const& q_pe,
torch::Tensor const& kv_c_and_k_pe_cache,
@@ -64,11 +65,11 @@ struct IsPersistent {
static const bool value = v;
};
-template >
+template >
struct MlaSm100 {
using Element = T;
using ElementAcc = float;
- using ElementOut = T;
+ using ElementOut = TOut;
using TileShape = Shape<_128, _128, Shape<_512, _64>>;
using TileShapeH = cute::tuple_element_t<0, TileShape>;
@@ -99,6 +100,7 @@ struct MlaSm100 {
template
typename T::Fmha::Arguments args_from_options(
at::Tensor const& out,
+ at::Tensor const& lse,
at::Tensor const& q_nope,
at::Tensor const& q_pe,
at::Tensor const& kv_c_and_k_pe_cache,
@@ -162,12 +164,15 @@ typename T::Fmha::Arguments args_from_options(
stride_PT,
page_count_total,
page_size},
- {static_cast(out.data_ptr()), stride_O, static_cast(nullptr), stride_LSE},
+ {static_cast(out.data_ptr()),
+ stride_O,
+ static_cast(lse.defined() ? lse.data_ptr() : nullptr),
+ stride_LSE},
hw_info,
// TODO(trevor-m): Change split_kv back to -1 when
// https://github.com/NVIDIA/cutlass/issues/2274 is fixed. Split_kv=1 will
// perform worse with larger context length and smaller batch sizes.
- num_kv_splits, // split_kv
+ static_cast(num_kv_splits), // split_kv
nullptr, // is_var_split_kv
};
// TODO(kaixih@nvidia): When split_kv=-1 and is_var_split_kv=false, we compute
@@ -178,9 +183,10 @@ typename T::Fmha::Arguments args_from_options(
return arguments;
}
-template
+template
void runMla(
at::Tensor const& out,
+ at::Tensor const& lse,
at::Tensor const& q_nope,
at::Tensor const& q_pe,
at::Tensor const& kv_c_and_k_pe_cache,
@@ -190,9 +196,9 @@ void runMla(
double sm_scale,
int64_t num_kv_splits,
cudaStream_t stream) {
- using MlaSm100Type = MlaSm100;
+ using MlaSm100Type = MlaSm100;
typename MlaSm100Type::Fmha fmha;
- auto arguments = args_from_options(out, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, sm_scale, num_kv_splits);
+ auto arguments = args_from_options(out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, sm_scale, num_kv_splits);
CUTLASS_CHECK(fmha.can_implement(arguments));
@@ -214,6 +220,7 @@ void runMla(
void sm100_cutlass_mla_decode(
torch::Tensor const& out,
+ torch::Tensor const& lse,
torch::Tensor const& q_nope,
torch::Tensor const& q_pe,
torch::Tensor const& kv_c_and_k_pe_cache,
@@ -233,14 +240,14 @@ void sm100_cutlass_mla_decode(
DISPATCH_BOOL(page_size == 128, IsPaged128, [&] {
DISPATCH_BOOL(num_kv_splits <= 1, NotManualSplitKV, [&] {
if (in_dtype == at::ScalarType::Half) {
- runMla>(
- out, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
+ runMla>(
+ out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
} else if (in_dtype == at::ScalarType::BFloat16) {
- runMla>(
- out, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
+ runMla>(
+ out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
} else if (in_dtype == at::ScalarType::Float8_e4m3fn) {
- runMla>(
- out, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
+ runMla>(
+ out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
} else {
TORCH_CHECK(false, "Unsupported input data type of MLA");
}
@@ -253,7 +260,7 @@ void sm100_cutlass_mla_decode(
int64_t sm100_cutlass_mla_get_workspace_size(int64_t max_seq_len, int64_t num_batches, int64_t sm_count, int64_t num_kv_splits) {
// Workspace size depends on ElementAcc and ElementLSE (same as ElementAcc)
// which are float, so Element type here doesn't matter.
- using MlaSm100Type = MlaSm100;
+ using MlaSm100Type = MlaSm100;
// Get split kv. Requires problem shape and sm_count only.
typename MlaSm100Type::Fmha::Arguments arguments;
@@ -264,7 +271,7 @@ int64_t sm100_cutlass_mla_get_workspace_size(int64_t max_seq_len, int64_t num_ba
// Assumes device 0 when getting sm_count.
arguments.hw_info.sm_count =
sm_count <= 0 ? cutlass::KernelHardwareInfo::query_device_multiprocessor_count(/*device_id=*/0) : sm_count;
- arguments.split_kv = num_kv_splits;
+ arguments.split_kv = static_cast(num_kv_splits);
MlaSm100Type::Fmha::set_split_kv(arguments);
return MlaSm100Type::Fmha::get_workspace_size(arguments);
diff --git a/csrc/cache.h b/csrc/cache.h
index 0970b704be..fd230bec27 100644
--- a/csrc/cache.h
+++ b/csrc/cache.h
@@ -40,9 +40,19 @@ void concat_and_cache_mla(torch::Tensor& kv_c, torch::Tensor& k_pe,
void convert_fp8(torch::Tensor& dst_cache, torch::Tensor& src_cache,
const double scale, const std::string& kv_cache_dtype);
-void gather_cache(
+void gather_and_maybe_dequant_cache(
torch::Tensor const& src_cache, // [NUM_BLOCKS, BLOCK_SIZE, ENTRIES...]
torch::Tensor const& dst, // [TOT_TOKENS, ENTRIES...]
torch::Tensor const& block_table, // [BATCH, BLOCK_INDICES]
torch::Tensor const& cu_seq_lens, // [BATCH+1]
- int64_t batch_size, std::optional seq_starts = std::nullopt);
\ No newline at end of file
+ int64_t batch_size, const std::string& kv_cache_dtype,
+ torch::Tensor const& scale,
+ std::optional seq_starts = std::nullopt);
+
+// TODO(hc): cp_gather_cache need support scaled kvcahe in the future.
+void cp_gather_cache(
+ torch::Tensor const& src_cache, // [NUM_BLOCKS, BLOCK_SIZE, ENTRIES...]
+ torch::Tensor const& dst, // [TOT_TOKENS, ENTRIES...]
+ torch::Tensor const& block_table, // [BATCH, BLOCK_INDICES]
+ torch::Tensor const& cu_seq_lens, // [BATCH+1]
+ int64_t batch_size, std::optional seq_starts = std::nullopt);
diff --git a/csrc/cache_kernels.cu b/csrc/cache_kernels.cu
index 131dcb15cd..80b4c47c55 100644
--- a/csrc/cache_kernels.cu
+++ b/csrc/cache_kernels.cu
@@ -1,6 +1,7 @@
#include
#include
#include
+#include
#include "cuda_utils.h"
#include "cuda_compat.h"
@@ -624,9 +625,9 @@ void convert_fp8(torch::Tensor& dst_cache, torch::Tensor& src_cache,
namespace vllm {
// grid is launched with dimensions (batch, num_splits)
-template
-__global__ void gather_cache(
- const scalar_t* __restrict__ src_cache, // [NUM_BLOCKS, BLOCK_SIZE,
+template
+__global__ void gather_and_maybe_dequant_cache(
+ const cache_t* __restrict__ src_cache, // [NUM_BLOCKS, BLOCK_SIZE,
// ENTRIES...]
scalar_t* __restrict__ dst, // [TOT_TOKENS, ENTRIES...]
const int32_t* __restrict__ block_table, // [BATCH, BLOCK_INDICES]
@@ -634,6 +635,7 @@ __global__ void gather_cache(
const int32_t block_size, const int32_t entry_size,
const int64_t block_table_stride, const int64_t cache_block_stride,
const int64_t cache_entry_stride, const int64_t dst_entry_stride,
+ const float* __restrict__ scale,
const int32_t* __restrict__ seq_starts) { // Optional: starting offsets per
// batch
@@ -675,10 +677,16 @@ __global__ void gather_cache(
if (partial_block_size) full_blocks_end -= 1;
}
- auto copy_entry = [&](const scalar_t* __restrict__ _src,
+ auto copy_entry = [&](const cache_t* __restrict__ _src,
scalar_t* __restrict__ _dst) {
- for (int i = threadIdx.x; i < entry_size; i += blockDim.x)
- _dst[i] = _src[i];
+ for (int i = threadIdx.x; i < entry_size; i += blockDim.x) {
+ if constexpr (kv_dt == Fp8KVCacheDataType::kAuto) {
+ _dst[i] = static_cast(_src[i]);
+ } else {
+ _dst[i] =
+ fp8::scaled_convert(_src[i], *scale);
+ }
+ }
};
for (int pid = split_start; pid < full_blocks_end; ++pid) {
@@ -705,8 +713,144 @@ __global__ void gather_cache(
} // namespace vllm
// Macro to dispatch the kernel based on the data type.
-#define CALL_GATHER_CACHE(CPY_DTYPE) \
- vllm::gather_cache<<>>( \
+// SCALAR_T is the data type of the destination tensor.
+// CACHE_T is the stored data type of kv-cache.
+// KV_DTYPE is the real data type of kv-cache.
+#define CALL_GATHER_CACHE(SCALAR_T, CACHE_T, KV_DTYPE) \
+ vllm::gather_and_maybe_dequant_cache \
+ <<>>( \
+ reinterpret_cast(src_cache.data_ptr()), \
+ reinterpret_cast(dst.data_ptr()), \
+ block_table.data_ptr(), cu_seq_lens.data_ptr(), \
+ block_size, entry_size, block_table_stride, cache_block_stride, \
+ cache_entry_stride, dst_entry_stride, \
+ reinterpret_cast(scale.data_ptr()), seq_starts_ptr);
+
+// Gather sequences from the cache into the destination tensor.
+// - cu_seq_lens contains the cumulative sequence lengths for each batch
+// - block_table contains the cache block indices for each sequence
+// - Optionally, seq_starts (if provided) offsets the starting block index by
+// (seq_starts[bid] / page_size)
+void gather_and_maybe_dequant_cache(
+ torch::Tensor const& src_cache, // [NUM_BLOCKS, BLOCK_SIZE, ENTRIES...]
+ torch::Tensor const& dst, // [TOT_TOKENS, ENTRIES...]
+ torch::Tensor const& block_table, // [BATCH, BLOCK_INDICES]
+ torch::Tensor const& cu_seq_lens, // [BATCH+1]
+ int64_t batch_size, const std::string& kv_cache_dtype,
+ torch::Tensor const& scale,
+ std::optional seq_starts = std::nullopt) {
+ at::cuda::OptionalCUDAGuard device_guard(src_cache.device());
+ const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+
+ int32_t block_size = src_cache.size(1);
+ int32_t entry_size = src_cache.flatten(2, -1).size(2);
+
+ TORCH_CHECK(block_table.dtype() == torch::kInt32,
+ "block_table must be int32");
+ TORCH_CHECK(cu_seq_lens.dtype() == torch::kInt32,
+ "cu_seq_lens must be int32");
+ if (seq_starts.has_value()) {
+ TORCH_CHECK(seq_starts.value().dtype() == torch::kInt32,
+ "seq_starts must be int32");
+ }
+
+ TORCH_CHECK(src_cache.device() == dst.device(),
+ "src_cache and dst must be on the same device");
+ TORCH_CHECK(src_cache.device() == block_table.device(),
+ "src_cache and block_table must be on the same device");
+ TORCH_CHECK(src_cache.device() == cu_seq_lens.device(),
+ "src_cache and cu_seq_lens must be on the same device");
+ if (seq_starts.has_value()) {
+ TORCH_CHECK(src_cache.device() == seq_starts.value().device(),
+ "src_cache and seq_starts must be on the same device");
+ }
+
+ int64_t block_table_stride = block_table.stride(0);
+ int64_t cache_block_stride = src_cache.stride(0);
+ int64_t cache_entry_stride = src_cache.stride(1);
+ int64_t dst_entry_stride = dst.stride(0);
+
+ // Decide on the number of splits based on the batch size.
+ int num_splits = batch_size > 128 ? 2 : batch_size > 64 ? 4 : 16;
+ dim3 grid(batch_size, num_splits);
+ dim3 block(1024);
+
+ const int32_t* seq_starts_ptr =
+ seq_starts.has_value() ? seq_starts.value().data_ptr() : nullptr;
+
+ DISPATCH_BY_KV_CACHE_DTYPE(dst.dtype(), kv_cache_dtype, CALL_GATHER_CACHE);
+}
+
+namespace vllm {
+template
+// Note(hc): The cp_gather_cache allows seq_starts to no longer be divisible by
+// block_size.
+__global__ void cp_gather_cache(
+ const scalar_t* __restrict__ src_cache, // [NUM_BLOCKS, BLOCK_SIZE,
+ // ENTRY_SIZE]
+ scalar_t* __restrict__ dst, // [TOT_TOKENS, ENTRY_SIZE]
+ const int32_t* __restrict__ block_table, // [BATCH, BLOCK_INDICES]
+ const int32_t* __restrict__ cu_seq_lens, // [BATCH+1]
+ const int32_t block_size, const int32_t entry_size,
+ const int64_t block_table_stride, const int64_t cache_block_stride,
+ const int64_t cache_entry_stride, const int64_t dst_entry_stride,
+ const int32_t* __restrict__ seq_starts // Optional: starting offsets per
+ // batch
+) {
+ const int64_t bid = blockIdx.x; // Batch ID
+ const int32_t num_splits = gridDim.y;
+ const int32_t split = blockIdx.y;
+ const int32_t seq_start = cu_seq_lens[bid];
+ const int32_t seq_end = cu_seq_lens[bid + 1];
+ const int32_t seq_len = seq_end - seq_start;
+ const int32_t tot_slots = seq_len;
+ const int32_t split_slots = cuda_utils::ceil_div(tot_slots, num_splits);
+
+ const int32_t split_start = split * split_slots;
+ const int32_t split_end = min((split + 1) * split_slots, tot_slots);
+
+ const bool is_active_split = (split_start < tot_slots);
+
+ if (!is_active_split) return;
+
+ // Adjust the pointer for the block_table for this batch.
+ // If seq_starts is provided, compute an offset based on it
+ const int32_t batch_offset = bid * block_table_stride;
+ int32_t offset = split_start;
+ if (seq_starts != nullptr) {
+ offset += seq_starts[bid];
+ }
+ int32_t offset_div = offset / block_size;
+ offset = offset % block_size;
+ const int32_t* batch_block_table = block_table + batch_offset;
+
+ // Adjust dst pointer based on the cumulative sequence lengths.
+ dst += seq_start * dst_entry_stride;
+
+ auto copy_entry = [&](const scalar_t* __restrict__ _src,
+ scalar_t* __restrict__ _dst) {
+ for (int i = threadIdx.x; i < entry_size; i += blockDim.x)
+ _dst[i] = _src[i];
+ };
+
+ for (int pid = split_start; pid < split_end; ++pid) {
+ auto block_id = batch_block_table[offset_div];
+ auto block_start_ptr = src_cache + block_id * cache_block_stride;
+ auto block_dst_ptr = dst + pid * dst_entry_stride;
+ copy_entry(block_start_ptr + offset * cache_entry_stride, block_dst_ptr);
+ offset += 1;
+ // bump to next block
+ if (offset == block_size) {
+ offset_div += 1;
+ offset = 0;
+ }
+ }
+}
+} // namespace vllm
+
+// Macro to dispatch the kernel based on the data type.
+#define CALL_CP_GATHER_CACHE(CPY_DTYPE) \
+ vllm::cp_gather_cache<<>>( \
reinterpret_cast(src_cache.data_ptr()), \
reinterpret_cast(dst.data_ptr()), \
block_table.data_ptr(), cu_seq_lens.data_ptr(), \
@@ -716,9 +860,9 @@ __global__ void gather_cache(
// Gather sequences from the cache into the destination tensor.
// - cu_seq_lens contains the cumulative sequence lengths for each batch
// - block_table contains the cache block indices for each sequence
-// - Optionally, seq_starts (if provided) offsets the starting block index by
-// (seq_starts[bid] / page_size)
-void gather_cache(
+// - Optionally, seq_starts (if provided) offsets the starting slot index by
+// seq_starts[bid]
+void cp_gather_cache(
torch::Tensor const& src_cache, // [NUM_BLOCKS, BLOCK_SIZE, ENTRIES...]
torch::Tensor const& dst, // [TOT_TOKENS, ENTRIES...]
torch::Tensor const& block_table, // [BATCH, BLOCK_INDICES]
@@ -769,11 +913,11 @@ void gather_cache(
seq_starts.has_value() ? seq_starts.value().data_ptr() : nullptr;
if (dtype_bits == 32) {
- CALL_GATHER_CACHE(uint32_t);
+ CALL_CP_GATHER_CACHE(uint32_t);
} else if (dtype_bits == 16) {
- CALL_GATHER_CACHE(uint16_t);
+ CALL_CP_GATHER_CACHE(uint16_t);
} else if (dtype_bits == 8) {
- CALL_GATHER_CACHE(uint8_t);
+ CALL_CP_GATHER_CACHE(uint8_t);
} else {
TORCH_CHECK(false, "Unsupported data type width: ", dtype_bits);
}
diff --git a/csrc/core/scalar_type.hpp b/csrc/core/scalar_type.hpp
index d0f85e2360..68a8750f58 100644
--- a/csrc/core/scalar_type.hpp
+++ b/csrc/core/scalar_type.hpp
@@ -321,6 +321,8 @@ static inline constexpr auto kFE3M2f =
ScalarType::float_(3, 2, true, ScalarType::NAN_NONE);
static inline constexpr auto kFE4M3fn =
ScalarType::float_(4, 3, true, ScalarType::NAN_EXTD_RANGE_MAX_MIN);
+static inline constexpr auto kFE8M0fnu =
+ ScalarType(8, 0, false, 0, true, ScalarType::NAN_EXTD_RANGE_MAX_MIN);
static inline constexpr auto kFE5M2 = ScalarType::float_IEEE754(5, 2);
static inline constexpr auto kFE8M7 = ScalarType::float_IEEE754(8, 7);
static inline constexpr auto kFE5M10 = ScalarType::float_IEEE754(5, 10);
diff --git a/csrc/cpu/cpu_types_x86.hpp b/csrc/cpu/cpu_types_x86.hpp
index 3952c43cbc..982f7c07a1 100644
--- a/csrc/cpu/cpu_types_x86.hpp
+++ b/csrc/cpu/cpu_types_x86.hpp
@@ -89,7 +89,7 @@ struct FP16Vec16 : public Vec {
explicit FP16Vec16(const FP32Vec16&);
- void save(void* ptr) const { *reinterpret_cast<__m256i*>(ptr) = reg; }
+ void save(void* ptr) const { _mm256_storeu_si256((__m256i*)ptr, reg); }
void save(void* ptr, const int elem_num) const {
constexpr uint32_t M = 0xFFFFFFFF;
@@ -126,7 +126,7 @@ struct BF16Vec16 : public Vec {
explicit BF16Vec16(const FP32Vec16&);
- void save(void* ptr) const { *reinterpret_cast<__m256i*>(ptr) = reg; }
+ void save(void* ptr) const { _mm256_storeu_si256((__m256i*)ptr, reg); }
void save(void* ptr, const int elem_num) const {
constexpr uint32_t M = 0xFFFFFFFF;
@@ -180,8 +180,8 @@ struct BF16Vec32 : public Vec {
(__m128i)vec8_data.reg, 1)) {}
void save(void* ptr) const {
- *reinterpret_cast<__m256i*>(ptr) = reg_low;
- *reinterpret_cast<__m256i*>((__m256i*)ptr + 1) = reg_high;
+ _mm256_storeu_si256((__m256i*)ptr, reg_low);
+ _mm256_storeu_si256((__m256i*)ptr + 1, reg_high);
}
};
#endif
diff --git a/csrc/cpu/dnnl_helper.cpp b/csrc/cpu/dnnl_helper.cpp
new file mode 100644
index 0000000000..6def0e061f
--- /dev/null
+++ b/csrc/cpu/dnnl_helper.cpp
@@ -0,0 +1,523 @@
+#include
+#include
+
+#include "common/memory_desc.hpp"
+#include "common/memory.hpp"
+
+#include "dnnl_helper.h"
+
+static dnnl::engine& default_engine() {
+ static dnnl::engine engine(dnnl::engine::kind::cpu, 0);
+ return engine;
+}
+
+static dnnl::stream& default_stream() {
+ static dnnl::stream stream(default_engine());
+ return stream;
+}
+
+void release_dnnl_matmul_handler(int64_t handler) {
+ DNNLMatMulPrimitiveHandler* ptr =
+ reinterpret_cast(handler);
+ delete ptr;
+}
+
+DNNLScratchPadManager::DNNLScratchPadManager() : size_(0), ptr_(nullptr) {
+ this->realloc(allocation_unit * 128);
+}
+
+void DNNLScratchPadManager::realloc(size_t new_size) {
+ new_size = round(new_size);
+ if (new_size > size_) {
+ ptr_ = std::aligned_alloc(64, new_size);
+ size_ = new_size;
+ }
+}
+
+DNNLScratchPadManager* DNNLScratchPadManager::get_dnnl_scratchpad_manager() {
+ static DNNLScratchPadManager manager;
+ return &manager;
+}
+
+template
+class DNNLPrimitiveCache {
+ public:
+ using cache_value_t = std::pair;
+ using result_value_t = VT;
+ using container_t = std::list;
+ using value_iterator_t = typename container_t::iterator;
+ using map_t = std::unordered_map;
+ using creator_t = VT (*)();
+
+ public:
+ DNNLPrimitiveCache(size_t capacity)
+ : capacity_(capacity),
+ values_(),
+ key_to_value_(std::min(256lu, capacity)) {
+ assert(capacity > 0);
+ }
+
+ template
+ result_value_t get_or_create(const KT& key, F&& creator) {
+ std::optional value = get_value(key);
+ if (value.has_value()) {
+ return value.value()->second;
+ } else {
+ return add_value({key, creator()})->second;
+ }
+ }
+
+ size_t size() const { return values_.size(); }
+
+ private:
+ void dump_data() {
+ std::stringstream ss;
+ ss << "table_id: " << std::hex << reinterpret_cast(this) << std::dec
+ << "\n";
+ ss << "container: [";
+ for (auto&& iter : values_) {
+ ss << "(" << iter.first << ", " << std::hex
+ << reinterpret_cast(iter.second.get()) << "), " << std::dec;
+ }
+ ss << "]\n";
+
+ ss << "map: [";
+ for (auto&& iter : key_to_value_) {
+ ss << "(" << iter.first << ", " << iter.second->first << ", " << std::hex
+ << reinterpret_cast(iter.second->second.get()) << std::dec
+ << "), ";
+ }
+ ss << "]\n";
+ std::printf("%s\n", ss.str().c_str());
+ }
+
+ value_iterator_t add_value(cache_value_t&& new_value) {
+ if (size() == capacity_) {
+ cache_value_t& last_item = values_.back();
+ key_to_value_.erase(last_item.first);
+ values_.pop_back();
+ }
+
+ auto& added_value_ = values_.emplace_front(std::move(new_value));
+ key_to_value_.emplace(added_value_.first, values_.begin());
+ return values_.begin();
+ }
+
+ std::optional get_value(const KT& key) {
+ if (key_to_value_.size() > 0 && key == values_.begin()->first) {
+ return values_.begin();
+ }
+
+ auto value_map_iterator = key_to_value_.find(key);
+ if (value_map_iterator != key_to_value_.end()) {
+ values_.splice(values_.begin(), values_, value_map_iterator->second);
+ return value_map_iterator->second;
+ } else {
+ return {};
+ }
+ }
+
+ private:
+ const size_t capacity_;
+ container_t values_;
+ map_t key_to_value_;
+};
+
+DNNLMatMulPrimitiveHandler::DNNLMatMulPrimitiveHandler(
+ const Args& args, dnnl::memory::data_type b_type)
+ : b_n_size_(args.b_n_size),
+ b_n_stride_(args.b_n_stride),
+ b_k_size_(args.b_k_size),
+ b_k_stride_(args.b_k_stride),
+ b_type_(b_type),
+ c_type_(args.c_type),
+ runtime_memory_ptrs_(8),
+ primitive_cache_size_(args.primitive_cache_size) {
+ assert(primitive_cache_size_ > 0);
+}
+
+void DNNLMatMulPrimitiveHandler::prepack_weight(
+ void* original_b_ptr, dnnl::memory::desc b_target_mem_desc) {
+ dnnl::memory::desc original_b_md({b_k_size_, b_n_size_}, b_type_,
+ {b_k_stride_, b_n_stride_});
+ dnnl::memory original_weight(original_b_md, default_engine(), original_b_ptr);
+ dnnl::memory packed_weight(b_target_mem_desc, default_engine());
+ {
+ dnnl::reorder(original_weight, packed_weight)
+ .execute(default_stream(), original_weight, packed_weight);
+ default_stream().wait();
+ }
+ memory_cache_[DNNL_ARG_WEIGHTS] = packed_weight;
+ b_target_mem_desc_ = b_target_mem_desc;
+}
+
+void DNNLMatMulPrimitiveHandler::set_runtime_memory_ptr(
+ size_t index, dnnl_memory* memory_ptr) {
+ dnnl::impl::memory_storage_t* mem_storage_ptr = memory_ptr->memory_storage();
+ dnnl_memory_desc* mem_desc = const_cast(memory_ptr->md());
+ runtime_memory_ptrs_[index] = {mem_storage_ptr, mem_desc};
+}
+
+std::pair
+DNNLMatMulPrimitiveHandler::get_runtime_memory_ptr(size_t index) {
+ return runtime_memory_ptrs_[index];
+}
+
+namespace std {
+template <>
+struct hash {
+ size_t operator()(
+ const W8A8MatMulPrimitiveHandler::ClassMatmulCacheKey& val) const {
+ return hash()(val.b_n_size) ^ hash()(val.b_k_size) ^
+ hash()(static_cast(val.a_qs)) ^
+ hash()(static_cast(val.b_qs)) ^ hash()(val.use_azp) ^
+ hash()(static_cast(val.c_type));
+ }
+};
+
+template <>
+struct hash {
+ size_t operator()(
+ const W8A8MatMulPrimitiveHandler::MSizeCacheKey& val) const {
+ return hash()(val.a_m_size) ^ hash()(val.use_bias) ^
+ hash()(static_cast(val.bias_type));
+ }
+};
+
+template <>
+struct hash {
+ size_t operator()(
+ const MatMulPrimitiveHandler::ClassMatmulCacheKey& val) const {
+ return hash()(val.b_n_size) ^ hash()(val.b_k_size);
+ }
+};
+
+template <>
+struct hash {
+ size_t operator()(const MatMulPrimitiveHandler::MSizeCacheKey& val) const {
+ return hash()(val.a_m_size) ^
+ hash()(val.a_m_stride) ^ hash()(val.use_bias) ^
+ hash()(static_cast(val.bias_type));
+ }
+};
+} // namespace std
+
+bool operator==(const W8A8MatMulPrimitiveHandler::ClassMatmulCacheKey& l,
+ const W8A8MatMulPrimitiveHandler::ClassMatmulCacheKey& r) {
+ return l.b_n_size == r.b_n_size && l.b_k_size == r.b_k_size &&
+ l.a_qs == r.a_qs && l.b_qs == r.b_qs && l.use_azp == r.use_azp &&
+ l.c_type == r.c_type;
+}
+
+bool operator==(const W8A8MatMulPrimitiveHandler::MSizeCacheKey& l,
+ const W8A8MatMulPrimitiveHandler::MSizeCacheKey& r) {
+ return l.use_bias == r.use_bias && l.a_m_size == r.a_m_size &&
+ l.bias_type == r.bias_type;
+}
+
+bool operator==(const MatMulPrimitiveHandler::ClassMatmulCacheKey& l,
+ const MatMulPrimitiveHandler::ClassMatmulCacheKey& r) {
+ return l.b_n_size == r.b_n_size && l.b_k_size == r.b_k_size;
+}
+
+bool operator==(const MatMulPrimitiveHandler::MSizeCacheKey& l,
+ const MatMulPrimitiveHandler::MSizeCacheKey& r) {
+ return l.a_m_size == r.a_m_size && l.a_m_stride == r.a_m_stride &&
+ l.use_bias == r.use_bias && l.bias_type == r.bias_type;
+}
+
+static std::shared_ptr
+get_w8a8_class_primitive_cache(
+ const W8A8MatMulPrimitiveHandler::ClassMatmulCacheKey& key,
+ int64_t cache_size) {
+ static W8A8MatMulPrimitiveHandler::ClassMatmulCache cache(128);
+ assert(cache_size > 0);
+ return cache.get_or_create(key, [&]() {
+ return std::make_shared(cache_size);
+ });
+}
+
+W8A8MatMulPrimitiveHandler::W8A8MatMulPrimitiveHandler(const Args& args)
+ : DNNLMatMulPrimitiveHandler(
+ static_cast(args),
+ dnnl::memory::data_type::s8),
+ use_azp_(args.use_a_zero_point),
+ a_qs_(args.a_quantization_strategy),
+ b_qs_(args.b_quantization_strategy),
+ m_size_cache_(nullptr) {
+ assert(a_qs_ != QuantizationStrategy::PER_OUTPUT_CHANNEL);
+ assert(b_qs_ != QuantizationStrategy::PER_TOKEN);
+ if (a_qs_ == QuantizationStrategy::PER_TOKEN) {
+ assert(!use_azp_);
+ };
+ prepack_weight(args.b_ptr,
+ create_primitive_desc(
+ MSizeCacheKey{.a_m_size = DNNL_RUNTIME_DIM_VAL,
+ .use_bias = false,
+ .bias_type = dnnl::memory::data_type::undef},
+ true)
+ .weights_desc());
+ init_runtime_memory_cache(args);
+}
+
+void W8A8MatMulPrimitiveHandler::execute(ExecArgs& args) {
+ auto&& [a_storage, a_mem_desc] = get_runtime_memory_ptr(0);
+ auto&& [c_storage, c_mem_desc] = get_runtime_memory_ptr(1);
+ a_storage->set_data_handle((void*)args.a_ptr);
+ a_mem_desc->dims[0] = args.a_m_size;
+ c_storage->set_data_handle((void*)args.c_ptr);
+ c_mem_desc->dims[0] = args.a_m_size;
+
+ if (a_qs_ == QuantizationStrategy::PER_TENSOR) {
+ auto&& [a_scale_storage, a_scale_mem_desc] = get_runtime_memory_ptr(2);
+ a_scale_storage->set_data_handle((void*)args.a_scales_ptr);
+ }
+ if (use_azp_) {
+ auto&& [a_zero_point_storage, a_zero_point_mem_desc] =
+ get_runtime_memory_ptr(3);
+ a_zero_point_storage->set_data_handle((void*)args.a_zero_points_ptr);
+ }
+
+ if (args.use_bias) {
+ auto&& [bias_storage, bias_mem_desc] = get_runtime_memory_ptr(4);
+ bias_storage->set_data_handle((void*)args.bias_ptr);
+ }
+
+ dnnl::matmul matmul = get_matmul_cache(args);
+
+ auto&& [scratchpad_storage, scratchpad_mem_desc] = get_runtime_memory_ptr(5);
+ scratchpad_storage->set_data_handle(
+ DNNLScratchPadManager::get_dnnl_scratchpad_manager()->get_data());
+
+ matmul.execute(default_stream(), memory_cache_);
+ default_stream().wait();
+}
+
+dnnl::matmul W8A8MatMulPrimitiveHandler::get_matmul_cache(
+ const MSizeCacheKey& key) {
+ if (m_size_cache_.get() == nullptr) {
+ ClassMatmulCacheKey key = {.b_n_size = b_n_size_,
+ .b_k_size = b_k_size_,
+ .a_qs = a_qs_,
+ .b_qs = b_qs_,
+ .use_azp = use_azp_,
+ .c_type = c_type_};
+ m_size_cache_ = get_w8a8_class_primitive_cache(key, primitive_cache_size_);
+ }
+
+ return m_size_cache_->get_or_create(key, [&]() {
+ dnnl::matmul::primitive_desc desc = this->create_primitive_desc(key, false);
+ auto manager = DNNLScratchPadManager::get_dnnl_scratchpad_manager();
+ manager->realloc(desc.scratchpad_desc().get_size());
+ return dnnl::matmul(desc);
+ });
+}
+
+void W8A8MatMulPrimitiveHandler::init_runtime_memory_cache(const Args& args) {
+ memory_cache_[DNNL_ARG_SRC] = dnnl::memory({{1, b_k_size_},
+ dnnl::memory::data_type::s8,
+ dnnl::memory::format_tag::ab},
+ default_engine(), nullptr);
+ set_runtime_memory_ptr(0, memory_cache_[DNNL_ARG_SRC].get());
+ memory_cache_[DNNL_ARG_DST] =
+ dnnl::memory({{1, b_n_size_}, c_type_, dnnl::memory::format_tag::ab},
+ default_engine(), nullptr);
+ set_runtime_memory_ptr(1, memory_cache_[DNNL_ARG_DST].get());
+
+ // For PER_TOKEN, scales will be applied in outside epilogue
+ if (a_qs_ == QuantizationStrategy::PER_TENSOR) {
+ memory_cache_[DNNL_ARG_ATTR_SCALES | DNNL_ARG_SRC] = dnnl::memory(
+ {{1}, dnnl::memory::data_type::f32, {1}}, default_engine(), nullptr);
+ set_runtime_memory_ptr(
+ 2, memory_cache_[DNNL_ARG_ATTR_SCALES | DNNL_ARG_SRC].get());
+ if (use_azp_) {
+ memory_cache_[DNNL_ARG_ATTR_ZERO_POINTS | DNNL_ARG_SRC] = dnnl::memory(
+ {{1}, dnnl::memory::data_type::s32, {1}}, default_engine(), nullptr);
+ set_runtime_memory_ptr(
+ 3, memory_cache_[DNNL_ARG_ATTR_ZERO_POINTS | DNNL_ARG_SRC].get());
+ }
+ }
+
+ if (b_qs_ == QuantizationStrategy::PER_TENSOR) {
+ memory_cache_[DNNL_ARG_ATTR_SCALES | DNNL_ARG_WEIGHTS] =
+ dnnl::memory({{1}, dnnl::memory::data_type::f32, {1}}, default_engine(),
+ (void*)args.b_scales_ptr);
+ } else if (b_qs_ == QuantizationStrategy::PER_OUTPUT_CHANNEL) {
+ memory_cache_[DNNL_ARG_ATTR_SCALES | DNNL_ARG_WEIGHTS] =
+ dnnl::memory({{b_n_size_}, dnnl::memory::data_type::f32, {1}},
+ default_engine(), (void*)args.b_scales_ptr);
+ }
+
+ memory_cache_[DNNL_ARG_BIAS] =
+ dnnl::memory({{b_n_size_}, dnnl::memory::data_type::f32, {1}},
+ default_engine(), nullptr);
+ set_runtime_memory_ptr(4, memory_cache_[DNNL_ARG_BIAS].get());
+
+ memory_cache_[DNNL_ARG_SCRATCHPAD] =
+ dnnl::memory({{b_n_size_}, dnnl::memory::data_type::f32, {1}},
+ default_engine(), nullptr);
+ set_runtime_memory_ptr(5, memory_cache_[DNNL_ARG_SCRATCHPAD].get());
+}
+
+dnnl::matmul::primitive_desc W8A8MatMulPrimitiveHandler::create_primitive_desc(
+ const MSizeCacheKey& key, bool first_time) {
+ dnnl::memory::desc a_md({key.a_m_size, b_k_size_},
+ dnnl::memory::data_type::s8,
+ dnnl::memory::format_tag::ab);
+ dnnl::memory::desc b_md;
+ if (first_time) {
+ b_md =
+ dnnl::memory::desc({b_k_size_, b_n_size_}, dnnl::memory::data_type::s8,
+ dnnl::memory::format_tag::any);
+ } else {
+ b_md = b_target_mem_desc_;
+ }
+ dnnl::memory::desc c_md({key.a_m_size, b_n_size_}, c_type_,
+ dnnl::memory::format_tag::ab);
+
+ dnnl::primitive_attr attr;
+
+ attr.set_scratchpad_mode(dnnl::scratchpad_mode::user);
+
+ // For PER_TOKEN, scales will be applied in outside epilogue
+ if (a_qs_ == QuantizationStrategy::PER_TENSOR) {
+ attr.set_scales_mask(DNNL_ARG_SRC, 0);
+ if (use_azp_) {
+ attr.set_zero_points_mask(DNNL_ARG_SRC, 0);
+ }
+ }
+
+ if (b_qs_ == QuantizationStrategy::PER_TENSOR) {
+ attr.set_scales_mask(DNNL_ARG_WEIGHTS, 0);
+ } else if (b_qs_ == QuantizationStrategy::PER_OUTPUT_CHANNEL) {
+ attr.set_scales_mask(DNNL_ARG_WEIGHTS, 2);
+ }
+
+ if (key.use_bias) {
+ // For PER_TOKEN, bias will be applied in epilogue
+ assert(a_qs_ == QuantizationStrategy::PER_TENSOR);
+ dnnl::memory::desc bias_md({1, b_n_size_}, key.bias_type, {b_n_size_, 1});
+ return dnnl::matmul::primitive_desc(default_engine(), a_md, b_md, bias_md,
+ c_md, attr);
+ } else {
+ return dnnl::matmul::primitive_desc(default_engine(), a_md, b_md, c_md,
+ attr);
+ }
+}
+
+MatMulPrimitiveHandler::MatMulPrimitiveHandler(const Args& args)
+ : DNNLMatMulPrimitiveHandler(
+ static_cast(args), args.ab_type),
+ m_size_cache_(nullptr) {
+ assert(ab_type_ == dnnl::memory::data_type::f32 ||
+ ab_type_ == dnnl::memory::data_type::bf16 ||
+ ab_type_ == dnnl::memory::data_type::f16);
+ prepack_weight(args.b_ptr,
+ create_primitive_desc(
+ MSizeCacheKey{.a_m_size = DNNL_RUNTIME_DIM_VAL,
+ .a_m_stride = DNNL_RUNTIME_DIM_VAL,
+ .use_bias = false,
+ .bias_type = dnnl::memory::data_type::undef},
+ true)
+ .weights_desc());
+ init_runtime_memory_cache(args);
+}
+
+static std::shared_ptr
+get_matul_class_primitive_cache(
+ const MatMulPrimitiveHandler::ClassMatmulCacheKey& key,
+ int64_t cache_size) {
+ static MatMulPrimitiveHandler::ClassMatmulCache cache(128);
+ assert(cache_size > 0);
+ return cache.get_or_create(key, [&]() {
+ return std::make_shared(cache_size);
+ });
+}
+
+void MatMulPrimitiveHandler::execute(ExecArgs& args) {
+ auto&& [a_storage, a_mem_desc] = get_runtime_memory_ptr(0);
+ auto&& [c_storage, c_mem_desc] = get_runtime_memory_ptr(1);
+ a_storage->set_data_handle((void*)args.a_ptr);
+ a_mem_desc->dims[0] = args.a_m_size;
+ a_mem_desc->format_desc.blocking.strides[0] = args.a_m_stride;
+ c_storage->set_data_handle((void*)args.c_ptr);
+ c_mem_desc->dims[0] = args.a_m_size;
+
+ if (args.use_bias) {
+ auto&& [bias_storage, bias_mem_desc] = get_runtime_memory_ptr(2);
+ bias_storage->set_data_handle((void*)args.bias_ptr);
+ }
+
+ dnnl::matmul matmul = get_matmul_cache(args);
+
+ auto&& [scratchpad_storage, scratchpad_mem_desc] = get_runtime_memory_ptr(3);
+ scratchpad_storage->set_data_handle(
+ DNNLScratchPadManager::get_dnnl_scratchpad_manager()->get_data());
+
+ matmul.execute(default_stream(), memory_cache_);
+ default_stream().wait();
+}
+
+dnnl::matmul MatMulPrimitiveHandler::get_matmul_cache(
+ const MSizeCacheKey& key) {
+ if (m_size_cache_.get() == nullptr) {
+ ClassMatmulCacheKey key = {.b_n_size = b_n_size_, .b_k_size = b_k_size_};
+ m_size_cache_ = get_matul_class_primitive_cache(key, primitive_cache_size_);
+ }
+ return m_size_cache_->get_or_create(key, [&]() {
+ dnnl::matmul::primitive_desc desc = this->create_primitive_desc(key, false);
+ auto manager = DNNLScratchPadManager::get_dnnl_scratchpad_manager();
+ manager->realloc(desc.scratchpad_desc().get_size());
+ return dnnl::matmul(desc);
+ });
+}
+
+dnnl::matmul::primitive_desc MatMulPrimitiveHandler::create_primitive_desc(
+ const MSizeCacheKey& key, bool first_time) {
+ dnnl::memory::desc a_md;
+ dnnl::memory::desc b_md;
+ if (first_time) {
+ a_md = dnnl::memory::desc({key.a_m_size, b_k_size_}, b_type_,
+ dnnl::memory::format_tag::ab);
+ b_md = dnnl::memory::desc({b_k_size_, b_n_size_}, b_type_,
+ dnnl::memory::format_tag::any);
+ } else {
+ a_md = dnnl::memory::desc({key.a_m_size, b_k_size_}, b_type_,
+ {key.a_m_stride, 1});
+ b_md = b_target_mem_desc_;
+ }
+ dnnl::memory::desc c_md({key.a_m_size, b_n_size_}, c_type_,
+ dnnl::memory::format_tag::ab);
+
+ dnnl::primitive_attr attr;
+ attr.set_scratchpad_mode(dnnl::scratchpad_mode::user);
+
+ if (key.use_bias) {
+ dnnl::memory::desc bias_md({1, b_n_size_}, key.bias_type, {b_n_size_, 1});
+ return dnnl::matmul::primitive_desc(default_engine(), a_md, b_md, bias_md,
+ c_md, attr);
+ } else {
+ return dnnl::matmul::primitive_desc(default_engine(), a_md, b_md, c_md,
+ attr);
+ }
+}
+
+void MatMulPrimitiveHandler::init_runtime_memory_cache(const Args& args) {
+ memory_cache_[DNNL_ARG_SRC] = dnnl::memory(
+ {{1, b_k_size_}, b_type_, {b_k_size_, 1}}, default_engine(), nullptr);
+ set_runtime_memory_ptr(0, memory_cache_[DNNL_ARG_SRC].get());
+ memory_cache_[DNNL_ARG_DST] =
+ dnnl::memory({{1, b_n_size_}, c_type_, dnnl::memory::format_tag::ab},
+ default_engine(), nullptr);
+ set_runtime_memory_ptr(1, memory_cache_[DNNL_ARG_DST].get());
+
+ memory_cache_[DNNL_ARG_BIAS] =
+ dnnl::memory({{b_n_size_}, dnnl::memory::data_type::f32, {1}},
+ default_engine(), nullptr);
+ set_runtime_memory_ptr(2, memory_cache_[DNNL_ARG_BIAS].get());
+
+ memory_cache_[DNNL_ARG_SCRATCHPAD] =
+ dnnl::memory({{b_n_size_}, dnnl::memory::data_type::f32, {1}},
+ default_engine(), nullptr);
+ set_runtime_memory_ptr(3, memory_cache_[DNNL_ARG_SCRATCHPAD].get());
+}
diff --git a/csrc/cpu/dnnl_helper.h b/csrc/cpu/dnnl_helper.h
new file mode 100644
index 0000000000..ad6773d2b9
--- /dev/null
+++ b/csrc/cpu/dnnl_helper.h
@@ -0,0 +1,243 @@
+#ifndef DNNL_HELPER_H
+#define DNNL_HELPER_H
+
+#include
+#include
+
+#include "oneapi/dnnl/dnnl.hpp"
+
+namespace c10 {
+struct BFloat16;
+struct Half;
+} // namespace c10
+
+namespace dnnl {
+namespace impl {
+struct memory_storage_t;
+struct matmul_pd_t;
+struct matmul_desc_t;
+} // namespace impl
+} // namespace dnnl
+struct dnnl_memory_desc;
+
+template
+class DNNLPrimitiveCache;
+
+template
+struct DNNLType {
+ static constexpr dnnl::memory::data_type type =
+ dnnl::memory::data_type::undef;
+};
+
+template <>
+struct DNNLType {
+ static constexpr dnnl::memory::data_type type = dnnl::memory::data_type::s8;
+};
+
+template <>
+struct DNNLType {
+ static constexpr dnnl::memory::data_type type = dnnl::memory::data_type::s32;
+};
+
+template <>
+struct DNNLType {
+ static constexpr dnnl::memory::data_type type = dnnl::memory::data_type::f32;
+};
+
+template <>
+struct DNNLType {
+ static constexpr dnnl::memory::data_type type = dnnl::memory::data_type::bf16;
+};
+
+template <>
+struct DNNLType {
+ static constexpr dnnl::memory::data_type type = dnnl::memory::data_type::f16;
+};
+
+template
+constexpr inline dnnl::memory::data_type get_dnnl_type() {
+ return DNNLType>::type;
+}
+
+class DNNLScratchPadManager {
+ public:
+ static constexpr size_t allocation_unit = 4 * 1024 * 1024; // 4KB
+
+ static DNNLScratchPadManager* get_dnnl_scratchpad_manager();
+
+ DNNLScratchPadManager();
+
+ template
+ T* get_data() {
+ return reinterpret_cast(ptr_);
+ }
+
+ static size_t round(size_t size) {
+ return ((size + allocation_unit - 1) / allocation_unit) * allocation_unit;
+ }
+
+ void realloc(size_t new_size);
+
+ private:
+ size_t size_;
+ void* ptr_;
+};
+
+class DNNLMatMulPrimitiveHandler {
+ public:
+ virtual ~DNNLMatMulPrimitiveHandler() = default;
+
+ protected:
+ struct Args {
+ dnnl_dim_t b_n_size;
+ dnnl_dim_t b_n_stride;
+ dnnl_dim_t b_k_size;
+ dnnl_dim_t b_k_stride;
+ void* b_ptr;
+ dnnl::memory::data_type c_type;
+ size_t primitive_cache_size;
+ };
+
+ protected:
+ DNNLMatMulPrimitiveHandler(const Args& args, dnnl::memory::data_type b_type);
+
+ void prepack_weight(void* original_b_ptr,
+ dnnl::memory::desc b_target_mem_desc);
+
+ void set_runtime_memory_ptr(size_t index, dnnl_memory* memory_ptr);
+
+ std::pair
+ get_runtime_memory_ptr(size_t index);
+
+ protected:
+ const dnnl_dim_t b_n_size_;
+ const dnnl_dim_t b_n_stride_;
+ const dnnl_dim_t b_k_size_;
+ const dnnl_dim_t b_k_stride_;
+ dnnl::memory::data_type b_type_;
+ dnnl::memory::data_type c_type_;
+ std::unordered_map memory_cache_;
+ std::vector>
+ runtime_memory_ptrs_;
+ dnnl::memory::desc b_target_mem_desc_;
+ int64_t primitive_cache_size_;
+};
+
+class W8A8MatMulPrimitiveHandler : public DNNLMatMulPrimitiveHandler {
+ public:
+ enum class QuantizationStrategy { PER_TOKEN, PER_TENSOR, PER_OUTPUT_CHANNEL };
+
+ struct Args : public DNNLMatMulPrimitiveHandler::Args {
+ bool use_a_zero_point;
+ QuantizationStrategy a_quantization_strategy;
+ QuantizationStrategy b_quantization_strategy;
+ float* b_scales_ptr;
+ };
+
+ struct ClassMatmulCacheKey {
+ dnnl_dim_t b_n_size;
+ dnnl_dim_t b_k_size;
+ QuantizationStrategy a_qs;
+ QuantizationStrategy b_qs;
+ bool use_azp;
+ dnnl::memory::data_type c_type;
+
+ friend bool operator==(const ClassMatmulCacheKey& l,
+ const ClassMatmulCacheKey& r);
+ };
+
+ struct MSizeCacheKey {
+ dnnl_dim_t a_m_size;
+ bool use_bias;
+ dnnl::memory::data_type bias_type;
+
+ friend bool operator==(const MSizeCacheKey& l, const MSizeCacheKey& r);
+ };
+
+ using MSizeCache = DNNLPrimitiveCache;
+ using ClassMatmulCache =
+ DNNLPrimitiveCache>;
+
+ struct ExecArgs : public MSizeCacheKey {
+ const int8_t* a_ptr;
+ const float* a_scales_ptr;
+ const int32_t* a_zero_points_ptr;
+ const void* bias_ptr;
+ void* c_ptr;
+ };
+
+ public:
+ W8A8MatMulPrimitiveHandler(const Args& args);
+
+ QuantizationStrategy get_input_scale_strategy() const { return a_qs_; }
+
+ bool get_input_use_zero_point() const { return use_azp_; }
+
+ void execute(ExecArgs& args);
+
+ private:
+ dnnl::matmul::primitive_desc create_primitive_desc(const MSizeCacheKey& key,
+ bool first_time);
+
+ void init_runtime_memory_cache(const Args& args);
+
+ dnnl::matmul get_matmul_cache(const MSizeCacheKey& key);
+
+ private:
+ const bool use_azp_;
+ const QuantizationStrategy a_qs_;
+ const QuantizationStrategy b_qs_;
+ std::shared_ptr m_size_cache_;
+};
+
+class MatMulPrimitiveHandler : public DNNLMatMulPrimitiveHandler {
+ public:
+ struct Args : public DNNLMatMulPrimitiveHandler::Args {
+ dnnl::memory::data_type ab_type;
+ };
+
+ struct ClassMatmulCacheKey {
+ dnnl_dim_t b_n_size;
+ dnnl_dim_t b_k_size;
+
+ friend bool operator==(const ClassMatmulCacheKey& l,
+ const ClassMatmulCacheKey& r);
+ };
+
+ struct MSizeCacheKey {
+ dnnl_dim_t a_m_size;
+ dnnl_dim_t a_m_stride;
+ bool use_bias;
+ dnnl::memory::data_type bias_type;
+
+ friend bool operator==(const MSizeCacheKey& l, const MSizeCacheKey& r);
+ };
+
+ using MSizeCache = DNNLPrimitiveCache;
+ using ClassMatmulCache =
+ DNNLPrimitiveCache>;
+
+ struct ExecArgs : public MSizeCacheKey {
+ const void* a_ptr;
+ const void* bias_ptr;
+ void* c_ptr;
+ };
+
+ public:
+ MatMulPrimitiveHandler(const Args& args);
+
+ void execute(ExecArgs& args);
+
+ private:
+ dnnl::matmul::primitive_desc create_primitive_desc(const MSizeCacheKey& key,
+ bool first_time);
+
+ void init_runtime_memory_cache(const Args& args);
+
+ dnnl::matmul get_matmul_cache(const MSizeCacheKey& key);
+
+ private:
+ std::shared_ptr m_size_cache_;
+};
+
+#endif
diff --git a/csrc/cpu/dnnl_helper.hpp b/csrc/cpu/dnnl_helper.hpp
deleted file mode 100644
index 1cb8dc5b25..0000000000
--- a/csrc/cpu/dnnl_helper.hpp
+++ /dev/null
@@ -1,206 +0,0 @@
-#ifndef DNNL_HELPER_HPP
-#define DNNL_HELPER_HPP
-
-#include
-#include
-
-#include "oneapi/dnnl/dnnl.hpp"
-
-namespace {
-template
-struct DNNLType {
- static constexpr dnnl::memory::data_type type =
- dnnl::memory::data_type::undef;
-};
-
-template <>
-struct DNNLType {
- static constexpr dnnl::memory::data_type type = dnnl::memory::data_type::s8;
-};
-
-template <>
-struct DNNLType {
- static constexpr dnnl::memory::data_type type = dnnl::memory::data_type::s32;
-};
-
-template <>
-struct DNNLType {
- static constexpr dnnl::memory::data_type type = dnnl::memory::data_type::f32;
-};
-
-template <>
-struct DNNLType {
- static constexpr dnnl::memory::data_type type = dnnl::memory::data_type::bf16;
-};
-
-template <>
-struct DNNLType {
- static constexpr dnnl::memory::data_type type = dnnl::memory::data_type::f16;
-};
-
-template
-constexpr inline dnnl::memory::data_type get_dnnl_type() {
- return DNNLType>::type;
-}
-}; // namespace
-
-template
-class DNNLPrimitiveHelper {
- public:
- // I8 input GEMM kernel (C = a_scales * A @ (b_scales * B^T) + bias)
- // A: [M, K], row-major
- // B: [K, N], column-major
- // C: [M, N], row-major
- // bias: [N], row-major, optional
- // a_scales: [MS]
- // b_scales: [NS]
- // Note: Due to the limitation of oneDNN
- // (https://github.com/oneapi-src/oneDNN/issues/1636), the quantized bias is
- // not supported.
-
- template
- static void gemm_s8s8_jit(const int8_t* a, const int8_t* b, OutputT* c,
- const BiasT* bias, dnnl_dim_t M, dnnl_dim_t N,
- dnnl_dim_t K, const float* a_scales,
- const float* b_scales, dnnl_dim_t MS,
- dnnl_dim_t NS) {
- auto&& OutputType = get_dnnl_type();
- auto&& BiasType = get_dnnl_type();
-
- dnnl::memory::desc a_md({M, K}, dnnl::memory::data_type::s8, {K, 1});
- dnnl::memory::desc b_md({K, N}, dnnl::memory::data_type::s8, {1, K});
- dnnl::memory::desc c_md({M, N}, OutputType, {N, 1});
-
- dnnl::primitive_attr attr;
- if constexpr (!InputNoScale) {
- if (MS == 1) {
- // per-tensor
- attr.set_scales_mask(DNNL_ARG_SRC, 0);
- } else {
- // per-token
- TORCH_CHECK(false, "per-token quantization is unsupported.");
- }
- }
-
- if (NS == 1) {
- // per-tensor
- attr.set_scales_mask(DNNL_ARG_WEIGHTS, 0);
- } else {
- // per-channel
- attr.set_scales_mask(DNNL_ARG_WEIGHTS, 2);
- }
-
- dnnl::matmul::primitive_desc matmul_pd;
-// Create memory descriptors with format_tag::any for the primitive. This
-// enables the matmul primitive to choose memory layouts for an
-// optimized primitive implementation, and these layouts may differ from the
-// ones provided by the user.
-#ifdef __aarch64__
- auto mat_src_md = dnnl::memory::desc({M, K}, dnnl::memory::data_type::s8,
- dnnl::memory::format_tag::any);
- auto mat_weights_md = dnnl::memory::desc(
- {K, N}, dnnl::memory::data_type::s8, dnnl::memory::format_tag::any);
- auto mat_dst_md =
- dnnl::memory::desc({M, N}, OutputType, dnnl::memory::format_tag::any);
- if (bias) {
- dnnl::memory::desc bias_md({1, N}, BiasType, {N, 1});
- matmul_pd = dnnl::matmul::primitive_desc(default_engine(), mat_src_md,
- mat_weights_md, bias_md,
- mat_dst_md, attr);
- } else {
- matmul_pd = dnnl::matmul::primitive_desc(
- default_engine(), mat_src_md, mat_weights_md, mat_dst_md, attr);
- }
-#else
- if (bias) {
- dnnl::memory::desc bias_md({1, N}, BiasType, {N, 1});
- matmul_pd = dnnl::matmul::primitive_desc(default_engine(), a_md, b_md,
- bias_md, c_md, attr);
- } else {
- matmul_pd = dnnl::matmul::primitive_desc(default_engine(), a_md, b_md,
- c_md, attr);
- }
-#endif
- dnnl::matmul matmul(matmul_pd);
-
- auto& engine = default_engine();
-
- dnnl::memory a_m(a_md, engine, (void*)a);
- dnnl::memory b_m(b_md, engine, (void*)b);
- dnnl::memory c_m(c_md, engine, (void*)c);
- dnnl::memory a_scales_m({{MS}, dnnl::memory::data_type::f32, {1}}, engine,
- (void*)a_scales);
- dnnl::memory b_scales_m({{NS}, dnnl::memory::data_type::f32, {1}}, engine,
- (void*)b_scales);
-
- auto& stream = default_stream();
-
- auto mat_src_mem = a_m;
- auto mat_weights_mem = b_m;
- auto mat_dst_mem = c_m;
-#ifdef __aarch64__
- if (matmul_pd.weights_desc() != b_m.get_desc()) {
- mat_weights_mem = dnnl::memory(matmul_pd.weights_desc(), engine);
- dnnl::reorder(b_m, mat_weights_mem).execute(stream, b_m, mat_weights_mem);
- }
-#endif
- if constexpr (InputNoScale) {
- if (bias) {
- dnnl::memory::desc bias_md({N}, BiasType, {1});
- dnnl::memory bias_m(bias_md, engine, (void*)bias);
- matmul.execute(
- stream, {
- {DNNL_ARG_SRC, mat_src_mem},
- {DNNL_ARG_WEIGHTS, mat_weights_mem},
- {DNNL_ARG_BIAS, bias_m},
- {DNNL_ARG_DST, mat_dst_mem},
- {DNNL_ARG_ATTR_SCALES | DNNL_ARG_WEIGHTS, b_scales_m},
- });
- } else {
- matmul.execute(
- stream, {
- {DNNL_ARG_SRC, mat_src_mem},
- {DNNL_ARG_WEIGHTS, mat_weights_mem},
- {DNNL_ARG_DST, mat_dst_mem},
- {DNNL_ARG_ATTR_SCALES | DNNL_ARG_WEIGHTS, b_scales_m},
- });
- }
- } else {
- if (bias) {
- dnnl::memory::desc bias_md({N}, BiasType, {1});
- dnnl::memory bias_m(bias_md, engine, (void*)bias);
- matmul.execute(
- stream, {
- {DNNL_ARG_SRC, mat_src_mem},
- {DNNL_ARG_WEIGHTS, mat_weights_mem},
- {DNNL_ARG_BIAS, bias_m},
- {DNNL_ARG_DST, mat_dst_mem},
- {DNNL_ARG_ATTR_SCALES | DNNL_ARG_SRC, a_scales_m},
- {DNNL_ARG_ATTR_SCALES | DNNL_ARG_WEIGHTS, b_scales_m},
- });
- } else {
- matmul.execute(
- stream, {
- {DNNL_ARG_SRC, mat_src_mem},
- {DNNL_ARG_WEIGHTS, mat_weights_mem},
- {DNNL_ARG_DST, mat_dst_mem},
- {DNNL_ARG_ATTR_SCALES | DNNL_ARG_SRC, a_scales_m},
- {DNNL_ARG_ATTR_SCALES | DNNL_ARG_WEIGHTS, b_scales_m},
- });
- }
- }
- stream.wait();
- }
-
- private:
- static dnnl::engine& default_engine() {
- static dnnl::engine engine(dnnl::engine::kind::cpu, 0);
- return engine;
- }
-
- static dnnl::stream& default_stream() {
- static dnnl::stream stream(default_engine());
- return stream;
- }
-};
-#endif
diff --git a/csrc/cpu/dnnl_kernels.cpp b/csrc/cpu/dnnl_kernels.cpp
new file mode 100644
index 0000000000..9a3af4ac9d
--- /dev/null
+++ b/csrc/cpu/dnnl_kernels.cpp
@@ -0,0 +1,549 @@
+#include "cpu_types.hpp"
+#include "dnnl_helper.h"
+
+namespace {
+template
+struct KernelVecType {
+ using load_vec_type = void;
+ using cvt_vec_type = void;
+};
+
+template <>
+struct KernelVecType {
+ using load_vec_type = vec_op::FP32Vec16;
+ using cvt_vec_type = vec_op::FP32Vec16;
+};
+
+#if !defined(__aarch64__) || defined(ARM_BF16_SUPPORT)
+template <>
+struct KernelVecType {
+ using load_vec_type = vec_op::BF16Vec16;
+ using cvt_vec_type = vec_op::FP32Vec16;
+};
+#endif
+
+template <>
+struct KernelVecType {
+#if defined(__powerpc64__) || defined(__s390x__)
+ // Power architecture-specific vector type
+ using load_vec_type = vec_op::FP32Vec16;
+#else
+ // Fallback for other architectures
+ using load_vec_type = vec_op::FP16Vec16;
+#endif
+ using cvt_vec_type = vec_op::FP32Vec16;
+};
+
+template
+void static_scaled_int8_quant_impl(const scalar_t* input, int8_t* output,
+ const float* scale, const int32_t* azp,
+ const int64_t num_tokens,
+ const int64_t input_stride,
+ const int64_t hidden_size) {
+ using load_vec_t = typename KernelVecType::load_vec_type;
+ using cvt_vec_t = typename KernelVecType::cvt_vec_type;
+ constexpr int64_t vec_elem_num = load_vec_t::VEC_ELEM_NUM;
+
+ constexpr float i8_min =
+ static_cast(std::numeric_limits::min());
+ constexpr float i8_max =
+ static_cast(std::numeric_limits::max());
+ const cvt_vec_t inv_scale(1.0 / *scale);
+ const cvt_vec_t i8_min_vec(i8_min);
+ const cvt_vec_t i8_max_vec(i8_max);
+
+ cvt_vec_t zp_vec;
+ if constexpr (AZP) {
+ zp_vec = cvt_vec_t(static_cast(*azp));
+ }
+
+#pragma omp parallel for
+ for (int64_t i = 0; i < num_tokens; ++i) {
+ int64_t j = 0;
+ const scalar_t* input_ptr = input + i * input_stride;
+ int8_t* output_ptr = output + i * hidden_size;
+ for (; j < hidden_size - vec_elem_num; j += vec_elem_num) {
+ load_vec_t elems(input_ptr + j);
+ cvt_vec_t elems_fp32(elems);
+ elems_fp32 = elems_fp32 * inv_scale;
+
+ if constexpr (AZP) {
+ elems_fp32 = elems_fp32 + zp_vec;
+ }
+
+ elems_fp32 = elems_fp32.clamp(i8_min_vec, i8_max_vec);
+ vec_op::INT8Vec16 elems_int8(elems_fp32);
+ elems_int8.save(output_ptr + j);
+ }
+
+ load_vec_t elems(input_ptr + j);
+ cvt_vec_t elems_fp32(elems);
+ elems_fp32 = elems_fp32 * inv_scale;
+
+ if constexpr (AZP) {
+ elems_fp32 = elems_fp32 + zp_vec;
+ }
+
+ elems_fp32 = elems_fp32.clamp(i8_min_vec, i8_max_vec);
+ vec_op::INT8Vec16 elems_int8(elems_fp32);
+ elems_int8.save(output_ptr + j, hidden_size - j);
+ }
+}
+
+template
+void dynamic_scaled_int8_quant_impl(const scalar_t* input, int8_t* output,
+ float* scale, int32_t* azp,
+ const int64_t num_tokens,
+ const int64_t input_stride,
+ const int64_t hidden_size) {
+ using load_vec_t = typename KernelVecType::load_vec_type;
+ using cvt_vec_t = typename KernelVecType::cvt_vec_type;
+ constexpr int vec_elem_num = load_vec_t::VEC_ELEM_NUM;
+
+ constexpr float i8_min =
+ static_cast(std::numeric_limits::min());
+ constexpr float i8_max =
+ static_cast(std::numeric_limits::max());
+ const cvt_vec_t i8_min_vec(i8_min);
+ const cvt_vec_t i8_max_vec(i8_max);
+
+#pragma omp parallel for
+ for (int64_t i = 0; i < num_tokens; ++i) {
+ cvt_vec_t max_value(std::numeric_limits::lowest());
+ cvt_vec_t min_value(std::numeric_limits::max());
+ {
+ int64_t j = 0;
+ const scalar_t* input_ptr = input + i * input_stride;
+ for (; j < hidden_size - vec_elem_num; j += vec_elem_num) {
+ load_vec_t elems(input_ptr + j);
+ cvt_vec_t elems_fp32(elems);
+ if constexpr (AZP) {
+ max_value = max_value.max(elems_fp32);
+ min_value = min_value.min(elems_fp32);
+ } else {
+ max_value = max_value.max(elems_fp32.abs());
+ }
+ }
+
+ load_vec_t elems(input_ptr + j);
+ cvt_vec_t elems_fp32(elems);
+
+ if (j + vec_elem_num == hidden_size) {
+ if constexpr (AZP) {
+ max_value = max_value.max(elems_fp32);
+ min_value = min_value.min(elems_fp32);
+ } else {
+ max_value = max_value.max(elems_fp32.abs());
+ }
+ } else {
+ if constexpr (AZP) {
+ max_value = max_value.max(elems_fp32, hidden_size - j);
+ min_value = min_value.min(elems_fp32, hidden_size - j);
+ } else {
+ max_value = max_value.max(elems_fp32.abs(), hidden_size - j);
+ }
+ }
+ }
+
+ float scale_val;
+ float azp_val = 0.0f;
+ if constexpr (AZP) {
+ float max_scalar = max_value.reduce_max();
+ float min_scalar = min_value.reduce_min();
+ scale_val = (max_scalar - min_scalar) / 255.0f;
+ azp_val = std::nearbyint(-128.0f - min_scalar / scale_val);
+ azp[i] = azp_val;
+ scale[i] = scale_val;
+ } else {
+ scale_val = max_value.reduce_max() / 127.0f;
+ scale[i] = scale_val;
+ }
+
+ const cvt_vec_t inv_scale(1.0 / scale_val);
+ const cvt_vec_t azp_vec(azp_val);
+
+ {
+ int64_t j = 0;
+ const scalar_t* input_ptr = input + i * input_stride;
+ int8_t* output_ptr = output + i * hidden_size;
+ for (; j < hidden_size - vec_elem_num; j += vec_elem_num) {
+ load_vec_t elems(input_ptr + j);
+ cvt_vec_t elems_fp32(elems);
+ elems_fp32 = (elems_fp32 * inv_scale);
+
+ if constexpr (AZP) {
+ elems_fp32 = elems_fp32 + azp_vec;
+ }
+ elems_fp32 = elems_fp32.clamp(i8_min_vec, i8_max_vec);
+ vec_op::INT8Vec16 elems_int8(elems_fp32);
+ elems_int8.save(output_ptr + j);
+ }
+
+ load_vec_t elems(input_ptr + j);
+ cvt_vec_t elems_fp32(elems);
+ elems_fp32 = (elems_fp32 * inv_scale);
+
+ if constexpr (AZP) {
+ elems_fp32 = elems_fp32 + azp_vec;
+ }
+ elems_fp32 = elems_fp32.clamp(i8_min_vec, i8_max_vec);
+ vec_op::INT8Vec16 elems_int8(elems_fp32);
+ elems_int8.save(output_ptr + j, hidden_size - j);
+ }
+ }
+}
+
+template
+void dynamic_quant_epilogue(const float* input, scalar_t* output,
+ const float* a_scale, const int32_t* azp,
+ const float* azp_adj, const scalar_t* bias,
+ const int64_t num_tokens,
+ const int64_t hidden_size) {
+ CPU_KERNEL_GUARD_IN(dynamic_quant_epilogue)
+ using load_vec_t = typename KernelVecType::load_vec_type;
+ using cvt_vec_t = typename KernelVecType::cvt_vec_type;
+ constexpr int vec_elem_num = load_vec_t::VEC_ELEM_NUM;
+
+ const int64_t thread_num = omp_get_max_threads();
+ if (num_tokens > thread_num) {
+#pragma omp parallel for
+ for (int64_t i = 0; i < num_tokens; ++i) {
+ const float* input_ptr = input + i * hidden_size;
+ scalar_t* output_ptr = output + i * hidden_size;
+ int64_t j = 0;
+ cvt_vec_t token_scale_vec(a_scale[i]);
+ cvt_vec_t token_zp_scale_vec;
+ if constexpr (AZP) {
+ float zp_scale_val = a_scale[i] * static_cast(azp[i]);
+ token_zp_scale_vec = cvt_vec_t(zp_scale_val);
+ }
+ for (; j < hidden_size - vec_elem_num; ++j) {
+ cvt_vec_t elems_fp32(input_ptr + j);
+ elems_fp32 = elems_fp32 * token_scale_vec;
+ if constexpr (AZP) {
+ cvt_vec_t azp_adj_fp32(azp_adj + j);
+ elems_fp32 = elems_fp32 - azp_adj_fp32 * token_zp_scale_vec;
+ }
+ if constexpr (Bias) {
+ load_vec_t bias_vec(bias + j);
+ cvt_vec_t bias_vec_fp32(bias_vec);
+ elems_fp32 = elems_fp32 + bias_vec_fp32;
+ }
+ load_vec_t elems_out(elems_fp32);
+ elems_out.save(output_ptr + j);
+ }
+ cvt_vec_t elems_fp32(input_ptr + j);
+ elems_fp32 = elems_fp32 * token_scale_vec;
+ if constexpr (AZP) {
+ cvt_vec_t azp_adj_fp32(azp_adj + j);
+ elems_fp32 = elems_fp32 - azp_adj_fp32 * token_zp_scale_vec;
+ }
+ if constexpr (Bias) {
+ load_vec_t bias_vec(bias + j);
+ cvt_vec_t bias_vec_fp32(bias_vec);
+ elems_fp32 = elems_fp32 + bias_vec_fp32;
+ }
+ load_vec_t elems_out(elems_fp32);
+ elems_out.save(output_ptr + j, hidden_size - j);
+ }
+ } else {
+ const int64_t vec_iteration =
+ (hidden_size + vec_elem_num - 1) / vec_elem_num;
+ const int64_t vec_iteration_per_thread =
+ (vec_iteration + thread_num - 1) / thread_num;
+ const int64_t elem_num_per_thread = vec_iteration_per_thread * vec_elem_num;
+#pragma omp parallel for schedule(static, 1)
+ for (int64_t i = 0; i < thread_num; ++i) {
+ const int64_t start = elem_num_per_thread * i;
+ const int64_t end = std::min(hidden_size, elem_num_per_thread + start);
+ for (int64_t j = 0; j < num_tokens; ++j) {
+ cvt_vec_t token_scale_vec(a_scale[j]);
+ cvt_vec_t token_zp_scale_vec;
+ if constexpr (AZP) {
+ float zp_scale_val = a_scale[j] * static_cast(azp[j]);
+ token_zp_scale_vec = cvt_vec_t(zp_scale_val);
+ }
+ int64_t k = start;
+ const float* input_ptr = input + j * hidden_size;
+ scalar_t* output_ptr = output + j * hidden_size;
+ for (; k < end - vec_elem_num; k += vec_elem_num) {
+ cvt_vec_t elems_fp32(input_ptr + k);
+ elems_fp32 = elems_fp32 * token_scale_vec;
+ if constexpr (AZP) {
+ cvt_vec_t azp_adj_fp32(azp_adj + k);
+ elems_fp32 = elems_fp32 - azp_adj_fp32 * token_zp_scale_vec;
+ }
+ if constexpr (Bias) {
+ load_vec_t bias_vec(bias + k);
+ cvt_vec_t bias_vec_fp32(bias_vec);
+ elems_fp32 = elems_fp32 + bias_vec_fp32;
+ }
+ load_vec_t elems_out(elems_fp32);
+ elems_out.save(output_ptr + k);
+ }
+ if (k < end) {
+ cvt_vec_t elems_fp32(input_ptr + k);
+ elems_fp32 = elems_fp32 * token_scale_vec;
+ if constexpr (AZP) {
+ cvt_vec_t azp_adj_fp32(azp_adj + k);
+ elems_fp32 = elems_fp32 - azp_adj_fp32 * token_zp_scale_vec;
+ }
+ if constexpr (Bias) {
+ load_vec_t bias_vec(bias + k);
+ cvt_vec_t bias_vec_fp32(bias_vec);
+ elems_fp32 = elems_fp32 + bias_vec_fp32;
+ }
+ load_vec_t elems_out(elems_fp32);
+ elems_out.save(output_ptr + k, end - k);
+ }
+ }
+ }
+ }
+}
+} // namespace
+
+int64_t create_onednn_scaled_mm_handler(
+ const torch::Tensor& b, // [IC, OC], column-major
+ const torch::Tensor& b_scales, // [1] or [OC]
+ at::ScalarType output_type, bool dynamic_act_quant, bool use_azp,
+ int64_t primitive_cache_size) {
+ TORCH_CHECK(b.dim() == 2);
+ TORCH_CHECK(b.stride(0) == 1); // Column-major
+ TORCH_CHECK(b_scales.is_contiguous());
+
+ W8A8MatMulPrimitiveHandler::Args args;
+ args.primitive_cache_size = primitive_cache_size;
+
+ if (b_scales.numel() == 1) {
+ args.b_quantization_strategy =
+ W8A8MatMulPrimitiveHandler::QuantizationStrategy::PER_TENSOR;
+ } else {
+ TORCH_CHECK_EQ(b_scales.numel(), b.size(1));
+ args.b_quantization_strategy =
+ W8A8MatMulPrimitiveHandler::QuantizationStrategy::PER_OUTPUT_CHANNEL;
+ }
+ args.b_scales_ptr = b_scales.data_ptr();
+ args.b_k_size = b.size(0);
+ args.b_k_stride = b.stride(0);
+ args.b_n_size = b.size(1);
+ args.b_n_stride = b.stride(1);
+ args.b_ptr = b.data_ptr();
+
+ if (dynamic_act_quant) {
+ // dynamic per-token, bias, A scales and A zps will be applied in outside.
+ args.a_quantization_strategy =
+ W8A8MatMulPrimitiveHandler::QuantizationStrategy::PER_TOKEN;
+ args.use_a_zero_point = false;
+ } else {
+ // static per-tensor
+ args.a_quantization_strategy =
+ W8A8MatMulPrimitiveHandler::QuantizationStrategy::PER_TENSOR;
+ args.use_a_zero_point = use_azp;
+ }
+
+ VLLM_DISPATCH_FLOATING_TYPES(output_type, "create_onednn_scaled_mm_handler",
+ [&] {
+ if (dynamic_act_quant) {
+ args.c_type = get_dnnl_type();
+ } else {
+ args.c_type = get_dnnl_type();
+ }
+ });
+
+ return reinterpret_cast(new W8A8MatMulPrimitiveHandler(args));
+}
+
+void onednn_scaled_mm(
+ torch::Tensor& c, // [M, OC], row-major
+ const torch::Tensor& a, // [M, IC], row-major
+ const torch::Tensor& a_scales, // [M] or [1]
+ const std::optional& azp, // [M] or [1]
+ const std::optional& azp_adj, // [M] or [1]
+ const std::optional& bias, // [N]
+ int64_t handler) {
+ CPU_KERNEL_GUARD_IN(onednn_scaled_mm)
+ TORCH_CHECK(a.dim() == 2);
+ TORCH_CHECK(a.is_contiguous());
+ TORCH_CHECK(c.is_contiguous());
+ W8A8MatMulPrimitiveHandler* ptr =
+ reinterpret_cast(handler);
+ const int32_t* azp_ptr = nullptr;
+ if (azp.has_value()) {
+ azp_ptr = azp->data_ptr();
+ }
+ if (ptr->get_input_scale_strategy() ==
+ W8A8MatMulPrimitiveHandler::QuantizationStrategy::PER_TENSOR) {
+ TORCH_CHECK_EQ(a_scales.numel(), 1);
+ }
+
+ W8A8MatMulPrimitiveHandler::ExecArgs exec_args;
+ exec_args.a_ptr = a.data_ptr();
+ exec_args.a_m_size = a.size(0);
+ exec_args.bias_ptr = nullptr;
+ exec_args.bias_type = get_dnnl_type();
+ exec_args.use_bias = false;
+ exec_args.a_scales_ptr = nullptr;
+ exec_args.a_zero_points_ptr = nullptr;
+
+ VLLM_DISPATCH_FLOATING_TYPES(c.scalar_type(), "onednn_scaled_mm", [&] {
+ if (ptr->get_input_scale_strategy() ==
+ W8A8MatMulPrimitiveHandler::QuantizationStrategy::PER_TENSOR) {
+ if (bias.has_value()) {
+ exec_args.bias_ptr = bias->data_ptr();
+ exec_args.bias_type = get_dnnl_type();
+ exec_args.use_bias = true;
+ }
+ exec_args.a_scales_ptr = a_scales.data_ptr();
+ exec_args.a_zero_points_ptr = azp_ptr;
+ exec_args.c_ptr = c.data_ptr();
+ ptr->execute(exec_args);
+ } else if (ptr->get_input_scale_strategy() ==
+ W8A8MatMulPrimitiveHandler::QuantizationStrategy::PER_TOKEN) {
+ torch::Tensor tmp_fp32_out =
+ torch::empty_like(c, ::at::ScalarType::Float);
+ exec_args.c_ptr = tmp_fp32_out.data_ptr();
+ ptr->execute(exec_args);
+ if (bias.has_value()) {
+ if (azp.has_value()) {
+ dynamic_quant_epilogue(
+ tmp_fp32_out.data_ptr(), c.data_ptr(),
+ a_scales.data_ptr(), azp_ptr, azp_adj->data_ptr(),
+ bias->data_ptr(), c.size(0), c.size(1));
+ } else {
+ dynamic_quant_epilogue(
+ tmp_fp32_out.data_ptr(), c.data_ptr(),
+ a_scales.data_ptr(), azp_ptr, nullptr,
+ bias->data_ptr(), c.size(0), c.size(1));
+ }
+ } else {
+ if (azp.has_value()) {
+ dynamic_quant_epilogue(
+ tmp_fp32_out.data_ptr(), c.data_ptr(),
+ a_scales.data_ptr(), azp_ptr, azp_adj->data_ptr(),
+ (scalar_t*)nullptr, c.size(0), c.size(1));
+ } else {
+ dynamic_quant_epilogue(
+ tmp_fp32_out.data_ptr(), c.data_ptr(),
+ a_scales.data_ptr(), azp_ptr, nullptr, (scalar_t*)nullptr,
+ c.size(0), c.size(1));
+ }
+ }
+ } else {
+ TORCH_CHECK(false, "invalid act quant type.");
+ }
+ });
+}
+
+// static-per-tensor quantization.
+void static_scaled_int8_quant(
+ torch::Tensor& out, // [batch, hidden_size]
+ const torch::Tensor& input, // [batch, hidden_size]
+ const torch::Tensor& scale, std::optional const& azp) {
+ CPU_KERNEL_GUARD_IN(static_scaled_int8_quant)
+ TORCH_CHECK(out.is_contiguous());
+ TORCH_CHECK_EQ(input.dim(), 2);
+ TORCH_CHECK_EQ(input.stride(1), 1);
+ TORCH_CHECK(scale.numel() == 1);
+ TORCH_CHECK(!azp.has_value() || azp->numel() == 1);
+
+ const int64_t stride = input.stride(0);
+ const int64_t hidden_size = input.size(1);
+ const int64_t num_tokens = input.size(0);
+ VLLM_DISPATCH_FLOATING_TYPES(
+ input.scalar_type(), "static_scaled_int8_quant_impl", [&] {
+ if (azp.has_value()) {
+ static_scaled_int8_quant_impl(
+ input.data_ptr(), out.data_ptr(),
+ scale.data_ptr(), azp->data_ptr(), num_tokens,
+ stride, hidden_size);
+ } else {
+ static_scaled_int8_quant_impl(input.data_ptr(),
+ out.data_ptr(),
+ scale.data_ptr(), nullptr,
+ num_tokens, stride, hidden_size);
+ }
+ });
+}
+
+// dynamic-per-token quantization.
+void dynamic_scaled_int8_quant(
+ torch::Tensor& out, // [batch, hidden_size]
+ const torch::Tensor& input, // [batch, hidden_size]
+ torch::Tensor& scale, // [batch, 1]
+ std::optional const& azp) {
+ CPU_KERNEL_GUARD_IN(dynamic_scaled_int8_quant)
+ TORCH_CHECK(out.is_contiguous());
+ TORCH_CHECK_EQ(input.dim(), 2);
+ TORCH_CHECK_EQ(input.stride(1), 1);
+
+ const int64_t hidden_size = input.size(1);
+ const int64_t num_tokens = input.size(0);
+ const int64_t stride = input.stride(0);
+ VLLM_DISPATCH_FLOATING_TYPES(
+ input.scalar_type(), "dynamic_scaled_int8_quant_impl", [&] {
+ if (azp.has_value()) {
+ dynamic_scaled_int8_quant_impl(
+ input.data_ptr(), out.data_ptr(),
+ scale.data_ptr(), azp->data_ptr(), num_tokens,
+ stride, hidden_size);
+ } else {
+ dynamic_scaled_int8_quant_impl(
+ input.data_ptr(), out.data_ptr(),
+ scale.data_ptr(), nullptr, num_tokens, stride,
+ hidden_size);
+ }
+ });
+}
+
+int64_t create_onednn_mm_handler(const torch::Tensor& b,
+ int64_t primitive_cache_size) {
+ TORCH_CHECK(b.dim() == 2);
+
+ MatMulPrimitiveHandler::Args args;
+ args.primitive_cache_size = primitive_cache_size;
+
+ args.b_k_size = b.size(0);
+ args.b_k_stride = b.stride(0);
+ args.b_n_size = b.size(1);
+ args.b_n_stride = b.stride(1);
+ args.b_ptr = b.data_ptr();
+
+ VLLM_DISPATCH_FLOATING_TYPES(b.scalar_type(), "create_onednn_mm_handler",
+ [&] {
+ args.c_type = get_dnnl_type();
+ args.ab_type = get_dnnl_type();
+ });
+
+ return reinterpret_cast(new MatMulPrimitiveHandler(args));
+}
+
+void onednn_mm(torch::Tensor& c, // [M, OC], row-major
+ const torch::Tensor& a, // [M, IC], row-major
+ const std::optional& bias, int64_t handler) {
+ CPU_KERNEL_GUARD_IN(onednn_mm)
+ TORCH_CHECK(a.dim() == 2);
+ TORCH_CHECK(a.stride(-1) == 1);
+ TORCH_CHECK(c.is_contiguous());
+ MatMulPrimitiveHandler* ptr =
+ reinterpret_cast(handler);
+
+ MatMulPrimitiveHandler::ExecArgs exec_args;
+ exec_args.a_m_size = a.size(0);
+ exec_args.a_m_stride = a.stride(0);
+
+ VLLM_DISPATCH_FLOATING_TYPES(a.scalar_type(), "onednn_mm", [&] {
+ if (bias.has_value()) {
+ exec_args.use_bias = true;
+ exec_args.bias_type = get_dnnl_type();
+ exec_args.bias_ptr = bias->data_ptr();
+ } else {
+ exec_args.use_bias = false;
+ exec_args.bias_type = get_dnnl_type();
+ exec_args.bias_ptr = nullptr;
+ }
+ exec_args.a_ptr = a.data_ptr();
+ exec_args.c_ptr = c.data_ptr();
+
+ ptr->execute(exec_args);
+ });
+}
diff --git a/csrc/cpu/quant.cpp b/csrc/cpu/quant.cpp
deleted file mode 100644
index 6e120b8d20..0000000000
--- a/csrc/cpu/quant.cpp
+++ /dev/null
@@ -1,951 +0,0 @@
-#include "cpu_types.hpp"
-#include "dnnl_helper.hpp"
-
-namespace {
-template
-struct KernelVecType {
- using load_vec_type = void;
- using azp_adj_load_vec_type = void;
- using cvt_vec_type = void;
-};
-
-template <>
-struct KernelVecType {
- using load_vec_type = vec_op::FP32Vec16;
- using azp_adj_load_vec_type = vec_op::INT32Vec16;
- using cvt_vec_type = vec_op::FP32Vec16;
-};
-
-#if !defined(__aarch64__) || defined(ARM_BF16_SUPPORT)
-template <>
-struct KernelVecType {
- using load_vec_type = vec_op::BF16Vec16;
- using azp_adj_load_vec_type = vec_op::INT32Vec16;
- using cvt_vec_type = vec_op::FP32Vec16;
-};
-#endif
-
-template <>
-struct KernelVecType {
-#if defined(__powerpc64__) || defined(__s390x__)
- // Power architecture-specific vector type
- using load_vec_type = vec_op::FP32Vec16;
-#else
- // Fallback for other architectures
- using load_vec_type = vec_op::FP16Vec16;
-#endif
- using azp_adj_load_vec_type = vec_op::INT32Vec16;
- using cvt_vec_type = vec_op::FP32Vec16;
-};
-
-#if defined(__AVX512F__) || defined(__aarch64__)
-template
-void static_scaled_int8_quant_impl(const scalar_t* input, int8_t* output,
- const float* scale, const int32_t* azp,
- const int num_tokens,
- const int hidden_size) {
- using load_vec_t = typename KernelVecType::load_vec_type;
- using cvt_vec_t = typename KernelVecType::cvt_vec_type;
- constexpr int vec_elem_num = load_vec_t::VEC_ELEM_NUM;
-
- constexpr float i8_min =
- static_cast(std::numeric_limits::min());
- constexpr float i8_max =
- static_cast(std::numeric_limits::max());
- const cvt_vec_t inv_scale(1.0 / *scale);
- const cvt_vec_t i8_min_vec(i8_min);
- const cvt_vec_t i8_max_vec(i8_max);
-
- cvt_vec_t zp_vec;
- if constexpr (AZP) {
- zp_vec = cvt_vec_t(static_cast